前言

就目前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];

}

}