前言
就目前hash算力而言，PC上的CPU和GPU是怎么也比不了专门的硬件的，所以本文只是出于学习的目的而写。
GPU编程概念简述
GPU编程，是将庞大的计算任务，分割成独立小任务后，交给GPU中动辄数
百上千甚至上万的处理单元，并行处理。比如要计算2000个 8*8 ,如果是cpu处理的话是要循
环2000次，但交给一个有2000个处理单元的GPU计算，每个处理单元计算一个8*8，则只需要一次即可全部计算完成。
关于Sha256原理
引用一篇文章，已把原理说的很明白了。
SHA256算法原理详解: https://blog.csdn.net/u011583927/article/details/80905740/
本文使用的sha256实现，是修改Bitcoin Core的sha256实现而来。
规则简述
1. sha256的规则是对要hash的数据体以512bit分割为一个块，并以前一个块为基础做64次逻辑计算。
2. 比特币挖矿是指对区块头做sha256 得到hash_1，然后再次对hash_1做sha256后得到hash_2
这个hash_2即是最终的区块hash，然后比对这个hash_2 如果<=全网难度，则挖矿成功。
实现思路
因只是用于比特币挖矿的hash(数据体长度固定)，而不是通用的hash(数据体长度不固定只能循环处理)
所以可以将区块头数据hash的指令全部展开为顺序执行，然后交给每个GPU的处理单元
这样，假如GPU有2000个处理单元，那么交给GPU一轮计算就可得到2000个hash结果。
实现细节
流程:
1. CPU端, 第一次补位: 是将比特币区块头数据补齐，如下:
区块头640bit(80byte) + 补位384bit(48byte) = 1024bit(512bit*2) = 128byte
2. 上述补位后的128byte数据交给GPU的每个处理单元，处理单元内的流程如下:
1. 替换nonce值，如下:
nonce在比特币头数据第76byte处，长度4个byte，将其修改为新的nonce值。
2. 将128byte分割为两个512bit的块，并按顺序对每个块执行sha256的64个逻辑运算得到结果hash_1。
3. 将hash_1补位，如下:
hash_1 32byte(256bit) + 补位32byte = 512bit = 64byte
4. 再次对补位后的hash_1做sha256运算，得到结果hash_2，此hash_2即是最终block hash.
测试结果
2.5 GHz 四核Intel Core i7
16 GB 1600 MHz DDR3
Intel Iris Pro 1536 MB
测试hash次数 20,480,000 次
CPU 消耗时间: 115.809秒
GPU 消耗时间: 11.477 秒
以下是部分代码摘录,全部代码请参考 比特币GPU独立挖矿节点中的以下文件
btc_sha256_gpu.h
btc_sha256_gpu.cpp
btc_sha256_gpu.cl
CPU端补位代码
// 将blockheader的80byte数据补位到128byte
void BTC_SHA256_GPU::PadBlockHaderData(TBlockHeaderBytes& vBytes) {
// 128 = 80(block头数据) + 48(sha256补位数据)
uint8_t* buf = &vBytes[0];
// 将后48byte置0
memset(buf+80, 0, 48);
// 补位
*(buf+80) = 0x80; // 数据末尾补 1000 0000
uint8_t rawDataSize[8];
WriteBE64(rawDataSize, 80 << 3); // 写入big endian (<<3的是乘以8)
memcpy(buf + 120, rawDataSize, 8);
}
GPU端Kernel部分代码
inline void DoHash(uint32_t* s, uint8_t* data, uint8_t blocks, uint8_t* out) {
Reset(s);
T(s, data, blocks);
Finalize(s, out);
}
/**
*
* @param pIn 根据sha256补位规则好的block头信息数组
* @param pOut 双sha256后的hash结果数组
* @param nonceOffset nonce的起始点
*/
__kernel void BtcDoubleSHA256(__global uint8_t* pIn, __global uint8_t* pOut, unsigned int nonceOffset) {
/**
* 双Sha256 Hash
* 1. pIn内是已经按照sha256规则补位好的128byte的区块头数据
* 128byte=1024bit=区块头(80byte=640bit)+补位(384bit=48byte)
* 2. 修改pIn内区块头内的nonce数据
* 3. 对修改后的pIn做一次sha256 hash运算
* 4. 对上述运算后的32byte的hash结果按sha256规则再次进行补位到512bit(64byte)
* 5. 再次对补位后的64byte做第二次sha256 hash运算
* 6. 将第二次hash结果输出
*/
uint32_t _s[8]; // sha256初始因子
uint8_t hashResult[64];
// 当前 puid
int puId = get_global_id(0);
// 当前需要处理的数据从global拷贝到private内存中
// private内存的读写速度优于global
uint8_t data[IN_DATA_SIZE];
#pragma unroll
for(int i =0; i < IN_DATA_SIZE; ++i) {
data[i] = pIn[i];
}
// 将nonce写入block header数据内
uint32_t nonce = puId + nonceOffset;
memcpy(data + NONCE_OFFSET, &nonce, 4);
// 第一次计算hash(128byte)
memset(hashResult, 0, 64);
DoHash(_s, data, 2, hashResult);
// 第二次hash前的补位(64byte)
hashResult[32] = 0x80; // sha256补位分隔符号
uint8_t rawDataSize[8];
WriteBE64_EX(rawDataSize, 32 << 3); // 写入big endian (<<3的是乘以8)
memcpy(hashResult + 56, rawDataSize, 8);
// 第二次计算hash
DoHash(_s, hashResult, 1, hashResult);
// 将本workitem的双hash结果拷贝到输出buf中
uint64_t outDataOffset = puId * OUT_DATA_SIZE;
#pragma unroll
for(int i =0; i < OUT_DATA_SIZE; ++i) {
pOut[i + outDataOffset] = hashResult[i];
}
}
