1.1编写目的本文的目的是分析CCminer开源代码中关于NeoScrypt算法的部分,并未涉及CCminer其他部分.文章主要围绕NeoScrypt算法的背景,算法流程,以及其在CCminer代码中是如何实现的,从而为后续NeoScrypt算法的优化工作做好铺垫. 1.2本文适用对象本文档是为需要研究NeoScrypt算法应用及其实现的人群提供的.考虑到本文的背景是POW,本文档的目标读者是了解POW机制的程序员. 1.3术语定义 Salt(盐):所谓加Salt，就是加点“佐料”。其基本想法是这样的——当用户首次提供密码时（通常是注册时），由系统自动往这个密码里撒一些“佐料”，然后再散列。而当用户登录时，系统为用户提供的代码撒上同样的“佐料”，然后散列，再比较散列值，已确定密码是否正确。　　第2章【算法介绍】 2.1背景相信对此文档感兴趣的读者对于虚拟货币和区块链已经有了一定的了解.莱特币和羽毛币所采用的算法就是我们今天要介绍的NeoScrypt算法.NeoScrypt算法的前身是Scrypt算法,正如名字所暗示的,NeoScrypt算法是对Scrypt算法的进一步发展. 2.2算法介绍前代Scrypt算法采用PBKDF2-HMAC-SHA256结构,用Salsa20/8算法将实际的数据混合到内存中,而NeoScrypt算法采用FastKDF-BLAKE2s结构,用Salsa20 / 20和ChaCha20 / 20将实际的数据混合到内存中.(备注:Salsa20/8与Salsa20 / 20是同一种算法,区别在于前者对数据进行了8轮处理,后者对数据进行了20轮处理). 2.3 Salsa20算法 Salsa20核函数可以将一个64字节的字节流x转换为另一个64字节的字节流Salsa20(x). 在此,我们将64字节的输入流x看作16个字 x0, x1, x2, ..., x15，其中每个字为4个字节的小端对齐的无符号数字.由于字只有4个字节，因此字的范围在{0,1,...,2^32-1}中.对这16个字进行总计320次可翻转的转换得到输出的16个字. 最后将输出的16个字分别加上原来的16个字并对2的32次方取模，同样保持小端字节序,从而得到64字节的输出字节流Salsa20(x). 每一个转换改变一个字z，先将两个其他的字x和y加和取2的32次方的模，再循环移位固定的距离后，将这个结果和z进行异或操作得到新的z.每次循环移位的距离是固定的.因此320次转换总的计算量包括320次加法、320次异或、320次循环移位操作. 整个转换过程包含10次相同的 double-round.每一个双倍循环又由两个rounds组成.每一个round由4个并行的1/4的round组成.每一个1/4的round转换4个字.整个转换过程对于每个字进行了20次的转换. 完整的函数定义如下: // R(a,b)宏定义了一个简单的循环左移位操作，rotl32(),C语言没有实现这个操作，汇编语言实现了 #define R(a,b) (((a) << (b)) | ((a) >> (32 - (b)))) void salsa20_word_specification(uint32 out[16],uint32 in[16]) { int i; uint32 x[16]; for (i = 0;i < 16;++i) x[i] = in[i]; // 这个循环10次 for (i = 20;i > 0;i -= 2) { // 循环移位的距离总是7,9,13,18，每4个语句为一组，例如前4个只对4,8,12,0进行转换 x[ 4] ^= R(x[ 0]+x[12], 7); x[ 8] ^= R(x[ 4]+x[ 0], 9); x[12] ^= R(x[ 8]+x[ 4],13); x[ 0] ^= R(x[12]+x[ 8],18); x[ 9] ^= R(x[ 5]+x[ 1], 7); x[13] ^= R(x[ 9]+x[ 5], 9); x[ 1] ^= R(x[13]+x[ 9],13); x[ 5] ^= R(x[ 1]+x[13],18); x[14] ^= R(x[10]+x[ 6], 7); x[ 2] ^= R(x[14]+x[10], 9); x[ 6] ^= R(x[ 2]+x[14],13); x[10] ^= R(x[ 6]+x[ 2],18); x[ 3] ^= R(x[15]+x[11], 7); x[ 7] ^= R(x[ 3]+x[15], 9); x[11] ^= R(x[ 7]+x[ 3],13); x[15] ^= R(x[11]+x[ 7],18); x[ 1] ^= R(x[ 0]+x[ 3], 7); x[ 2] ^= R(x[ 1]+x[ 0], 9); x[ 3] ^= R(x[ 2]+x[ 1],13); x[ 0] ^= R(x[ 3]+x[ 2],18); x[ 6] ^= R(x[ 5]+x[ 4], 7); x[ 7] ^= R(x[ 6]+x[ 5], 9); x[ 4] ^= R(x[ 7]+x[ 6],13); x[ 5] ^= R(x[ 4]+x[ 7],18); x[11] ^= R(x[10]+x[ 9], 7); x[ 8] ^= R(x[11]+x[10], 9); x[ 9] ^= R(x[ 8]+x[11],13); x[10] ^= R(x[ 9]+x[ 8],18); x[12] ^= R(x[15]+x[14], 7); x[13] ^= R(x[12]+x[15], 9); x[14] ^= R(x[13]+x[12],13); x[15] ^= R(x[14]+x[13],18); } // 最后转换后的字与原始字进行加和 for (i = 0;i < 16;++i) out[i] = x[i] + in[i]; } 2.4 ChaCha算法 ChaCha算法与Salsa算法基本相同,区别之处在于每次循环位移的距离不同.Salsa算法的循环位移距离是7,9,13,18, ChaCha算法的循环位移距离是16,12,8,7. 第3章代码实现 3.1CPU部分 ![](https://steemitimages.com/DQmWE2oS26A1bQgiPfKbA3JcXF7fMGXeZYsPT2wqoP7jZW9/image.png) neoscrypt((uchar*)vhash, (uchar*) endiandata, 0x80000620U); 参数: Vhash ---- 传入时为空,函数执行完毕后,传出一个处理后的Vhash Endiandata ---- 待处理数据 0x80000620U ---- 辅助作用,函数内部的条件判断将会用到从图中我们可以看出,NeoScrypt算法在CPU代码中主要分为四个部分: neoscrypt_fastkdf(input, 80, input, 80, 32, (uchar *) X, 256); 第一次大循环 第二次大循环 neoscrypt_fastkdf(input, 80, (uchar *) X, 256, 32, output, 32); ![](https://steemitimages.com/DQmU4ctM1wFZFbzCvNXmhc9Ej79f8zsPHWxdcJzFPmTucfN/image.png) neoscrypt_fastkdf(input, 80, input, 80, 32, (uchar *) X, 256); 参数: 第一个Input,80 ---- 80字节的key输入第二个Input,80 ---- 80字节的salt输入 X,256 ---- 256字节的output输出 ![](https://steemitimages.com/DQmPnQLCHAzgoNeUuvdMCamqcKQQUDjVEeZgsA1rt2AHSHH/image.png) ![](https://steemitimages.com/DQmUS91kBXRSr1UBUvsRuaSjJx7VB92ZNRvFcD8hBvCgsqR/image.png) 第一次和第二次大循环组成结构大致相同,区别在于对于数据处理的算法不同. 3.2GPU(CUDA)部分 ![](https://steemitimages.com/DQmVMwqNntFAXi78uuZYhAnK3LrUDuZXba8XPHQpPD1uJhi/image.png) 从图中可以看出,NeoScrypt算法的实现在GPU代码中主要分为6个部分: CudaMemset() ---- 内存设置 neoscrypt_gpu_hash_start <<>> () ---- 为该函数开启grid2*block2数量的线程 neoscrypt_gpu_hash_salsa1 <<>> () neoscrypt_gpu_hash_chacha1 <<>> () neoscrypt_gpu_hash_ending <<>> () cudaMemcpy(resNonces, d_NNonce[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost) ---- 将处理过的数据从GPU拷贝到CPU ![](https://steemitimages.com/DQmTrHrsEY1qHLCyHdHARuSfErXQP5KbRK5CXdQHbSeDD8W/image.png) ![](https://steemitimages.com/DQmX8FudUwakSnneJV2c6UpWr4PrJywBoNjNWbqrQPN8UGr/image.png) ![](https://steemitimages.com/DQmeqqcv3h48wrxaQnoRYKW6GNHCSuKKGKFVAQ3FPpaEwFM/image.png) ![](https://steemitimages.com/DQmfK46r2GrHzMbLmWVqAxxeptSbzdYbyu3jj7ixCXymwmV/image.png) 第4章结论对比NeoScrypt算法在CPU和GPU上的代码实现,不难看出,两者的结构构成很相似,大体流程都是fasktkdf → 两次大循环 → fastkdf. 追溯到底层,都是用Salsa算法或ChaCha算法对数据处理. 第5章参考文献 NeoScrypt算法白皮书 , http://phoenixcoin.org/archive/neoscrypt_v1.pdf Scrypt基于密码的密钥派生函数 , http://rossihwang.farbox.com/post/2014-03-25 Salsa20笔记 , http://www.mamicode.com/info-detail-2188736.html

lichengxin