#include
#include
#include "../include/utils.cuh"
#include
#include
#include
// TODO: Implement function to search for all nonces from 1 through MAX_NONCE (inclusive) using CUDA Threads
//extern BYTE difficulty_5_zeros[SHA256_HASH_SIZE];
//extern void printResult(BYTE *block_hash, uint64_t nonce, float seconds);
BYTE *d_block_content, *d_block_hash;
uint64_t *d_nonce;
__global__ void findNonce(BYTE* block_content, size_t current_length, BYTE* block_hash, uint64_t* nonce) {
BYTE difficulty[SHA256_HASH_SIZE] = "0000000000000000000000000000000000000000000000000000000000000000";
int index = blockIdx.x * blockDim.x + threadIdx.x;
BYTE c_block_content[BLOCK_SIZE + 30], c_block_hash[SHA256_HASH_SIZE];
d_strcpy((char *)c_block_content, (const char *)block_content);
if (index <= MAX_NONCE) {
char nonce_string[NONCE_SIZE];
intToString(index, nonce_string);
d_strcpy((char*)c_block_content + current_length, nonce_string);
apply_sha256(c_block_content, current_length + d_strlen(nonce_string), c_block_hash, 1);
if (compare_hashes(c_block_hash, difficulty) <= 0){
//printf("Block hash found: %s, below difficulty %s, for nonce %" PRIu64 "\n", block_hash, difficulty_5_zeros, *nonce);
//printResult(block_hash, *nonce, 0.0);
//(*nonce)++;
*nonce = index;
d_strcpy((char *)block_hash, (const char *)c_block_hash);
return;
}
}
}
int main(int argc, char **argv) {
BYTE hashed_tx1[SHA256_HASH_SIZE], hashed_tx2[SHA256_HASH_SIZE], hashed_tx3[SHA256_HASH_SIZE], hashed_tx4[SHA256_HASH_SIZE],
tx12[SHA256_HASH_SIZE * 2], tx34[SHA256_HASH_SIZE * 2], hashed_tx12[SHA256_HASH_SIZE], hashed_tx34[SHA256_HASH_SIZE],
tx1234[SHA256_HASH_SIZE * 2], top_hash[SHA256_HASH_SIZE], block_content[BLOCK_SIZE];
BYTE block_hash[SHA256_HASH_SIZE]; // TODO: Update
uint64_t nonce = 0; // TODO: Update
size_t current_length;
cudaMalloc((void**)&d_block_content, BLOCK_SIZE);
cudaMalloc((void**)&d_block_hash, SHA256_HASH_SIZE);
cudaMalloc((void**)&d_nonce, sizeof(uint64_t));
// Top hash
apply_sha256(tx1, strlen((const char*)tx1), hashed_tx1, 1);
apply_sha256(tx2, strlen((const char*)tx2), hashed_tx2, 1);
apply_sha256(tx3, strlen((const char*)tx3), hashed_tx3, 1);
apply_sha256(tx4, strlen((const char*)tx4), hashed_tx4, 1);
strcpy((char *)tx12, (const char *)hashed_tx1);
strcat((char *)tx12, (const char *)hashed_tx2);
apply_sha256(tx12, strlen((const char*)tx12), hashed_tx12, 1);
strcpy((char *)tx34, (const char *)hashed_tx3);
strcat((char *)tx34, (const char *)hashed_tx4);
apply_sha256(tx34, strlen((const char*)tx34), hashed_tx34, 1);
strcpy((char *)tx1234, (const char *)hashed_tx12);
strcat((char *)tx1234, (const char *)hashed_tx34);
apply_sha256(tx1234, strlen((const char*)tx34), top_hash, 1);
// prev_block_hash + top_hash
strcpy((char*)block_content, (const char*)prev_block_hash);
strcat((char*)block_content, (const char*)top_hash);
current_length = strlen((char*) block_content);
cudaEvent_t start, stop;
startTiming(&start, &stop);
cudaMemcpy(d_block_hash, block_hash, SHA256_HASH_SIZE, cudaMemcpyHostToDevice);
cudaMemcpy(d_block_content, block_content, BLOCK_SIZE, cudaMemcpyHostToDevice);
cudaMemcpy(d_nonce, &nonce, sizeof(uint64_t), cudaMemcpyHostToDevice);
const size_t block_size = 300;
size_t blocks_no = 100000000 / block_size;
if (100000000 % block_size)
++blocks_no;
findNonce<<>>(d_block_content, current_length, d_block_hash, d_nonce);
cudaDeviceSynchronize();
float seconds = stopTiming(&start, &stop);
cudaMemcpy(block_hash, d_block_hash, SHA256_HASH_SIZE, cudaMemcpyDeviceToHost);
cudaMemcpy(&nonce, d_nonce, sizeof(uint64_t), cudaMemcpyDeviceToHost);
cudaFree(d_block_content);
cudaFree(d_block_hash);
cudaFree(d_nonce);
//float seconds = stopTiming(&start, &stop);
printResult(block_hash, nonce, seconds);
return 0;