CCCPaste

e2b22b263dd73

#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;