From 9efe0b965d47de8e5449a7b9c2026907aec5266a Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Fri, 12 Sep 2014 23:28:19 +0200 Subject: [PATCH] blake: only use high part of target on gpu Add another few MH/s boost :) --- blake32.cu | 71 +++++++++++++++++++++++------------------------------- miner.h | 4 +-- 2 files changed, 32 insertions(+), 43 deletions(-) diff --git a/blake32.cu b/blake32.cu index 1ae994d..fd2c7c5 100644 --- a/blake32.cu +++ b/blake32.cu @@ -21,7 +21,7 @@ extern "C" uint32_t crc32_u32t(const uint32_t *buf, size_t size); extern "C" int blake256_rounds = 14; /* hash by cpu with blake 256 */ -extern "C" void blake256hash(void *output, const void *input, int rounds = 14) +extern "C" void blake256hash(void *output, const void *input, int8_t rounds = 14) { unsigned char hash[64]; sph_blake256_context ctx; @@ -44,15 +44,16 @@ extern bool opt_n_threads; extern bool opt_benchmark; extern int device_map[8]; -__constant__ -static uint32_t __align__(32) c_Target[8]; - __constant__ static uint32_t __align__(32) c_data[20]; +// only store the 2 high uint32 of the target hash +__constant__ static uint64_t c_Target; +__constant__ static int8_t c_BlakeRounds; + /* 8 adapters max (-t threads) */ -static uint32_t *d_resNounce[8]; -static uint32_t *h_resNounce[8]; +static uint32_t *d_resNonce[8]; +static uint32_t *h_resNonce[8]; /* max count of found nounces in one call */ #define NBN 2 @@ -131,7 +132,7 @@ static const uint32_t __align__(32) c_Padding[16] = { }; __device__ static -void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, int blakerounds) +void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0) { uint32_t /* __align__(8) */ m[16]; uint32_t /* __align__(8) */ v[16]; @@ -159,7 +160,8 @@ void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, in v[14] = c_u256[6]; v[15] = c_u256[7]; - for (int i = 0; i < blakerounds; i++) { + int rounds = c_BlakeRounds; + for (int i = 0; i < rounds; i++) { /* column step */ GS(0, 4, 0x8, 0xC, 0x0); GS(1, 5, 0x9, 0xD, 0x2); @@ -180,7 +182,7 @@ void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, in } __global__ -void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resNounce, const int blakerounds, const int crcsum) +void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resNounce, const int crcsum) { uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); if (thread < threads) @@ -194,11 +196,11 @@ void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resN } #if !USE_CACHE - blake256_compress(h, c_data, 512, blakerounds); + blake256_compress(h, c_data, 512); #else if (crcsum != prevsum) { prevsum = crcsum; - blake256_compress(h, c_data, 512, blakerounds); + blake256_compress(h, c_data, 512); #pragma unroll for(int i=0; i<8; i++) { cache[i] = h[i]; @@ -218,21 +220,10 @@ void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resN ending[2] = c_data[18]; ending[3] = nounce; /* our tested value */ - blake256_compress(h, ending, 640, blakerounds); -#if 0 - for (int i = 7; i >= 0; i--) { - uint32_t hash = cuda_swab32(h[i]); - if (hash > c_Target[i]) { - return; - } - if (hash < c_Target[i]) { - break; - } - } -#else + blake256_compress(h, ending, 640); + /* do not test all parts, fulltest() will do it */ - if (cuda_swab32(h[7]) <= c_Target[7]) -#endif + if (((uint64_t*)h)[3] <= c_Target) #if NBN == 2 /* keep the smallest nounce, + extra one if found */ if (resNounce[0] > nounce) { @@ -248,7 +239,7 @@ void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resN } __host__ -uint32_t blake256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, const int blakerounds, const uint32_t crcsum) +uint32_t blake256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, const uint32_t crcsum) { const int threadsperblock = TPB; uint32_t result = MAXU; @@ -258,32 +249,33 @@ uint32_t blake256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce size_t shared_size = 0; /* Check error on Ctrl+C or kill to prevent segfaults on exit */ - if (cudaMemset(d_resNounce[thr_id], 0xff, NBN*sizeof(uint32_t)) != cudaSuccess) + if (cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t)) != cudaSuccess) return result; - blake256_gpu_hash_80<<>>(threads, startNounce, d_resNounce[thr_id], blakerounds, crcsum); + blake256_gpu_hash_80<<>>(threads, startNounce, d_resNonce[thr_id], crcsum); cudaDeviceSynchronize(); - if (cudaSuccess == cudaMemcpy(h_resNounce[thr_id], d_resNounce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) { + if (cudaSuccess == cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) { //cudaThreadSynchronize(); /* seems no more required */ - result = h_resNounce[thr_id][0]; + result = h_resNonce[thr_id][0]; for (int n=0; n < (NBN-1); n++) - extra_results[n] = h_resNounce[thr_id][n+1]; + extra_results[n] = h_resNonce[thr_id][n+1]; } return result; } __host__ -void blake256_cpu_setBlock_80(uint32_t *pdata, const uint32_t *ptarget) +void blake256_cpu_setBlock_80(uint32_t *pdata, const uint32_t *ptarget, int8_t blakerounds) { uint32_t data[20]; memcpy(data, pdata, 80); CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_data, data, sizeof(data), 0, cudaMemcpyHostToDevice)); CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_sigma, host_sigma, sizeof(host_sigma), 0, cudaMemcpyHostToDevice)); - CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_Target, ptarget, 32, 0, cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_Target, &ptarget[6], 2*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_BlakeRounds, &blakerounds, sizeof(int8_t), 0, cudaMemcpyHostToDevice)); } extern "C" int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *ptarget, - uint32_t max_nonce, unsigned long *hashes_done, uint32_t blakerounds=14) + uint32_t max_nonce, unsigned long *hashes_done, int8_t blakerounds=14) { const uint32_t first_nonce = pdata[19]; static bool init[8] = { 0, 0, 0, 0, 0, 0, 0, 0 }; @@ -311,22 +303,19 @@ extern "C" int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *pt if (opt_n_threads > 1) { CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); } - CUDA_SAFE_CALL(cudaMallocHost(&h_resNounce[thr_id], NBN * sizeof(uint32_t))); - CUDA_SAFE_CALL(cudaMalloc(&d_resNounce[thr_id], NBN * sizeof(uint32_t))); + CUDA_SAFE_CALL(cudaMallocHost(&h_resNonce[thr_id], NBN * sizeof(uint32_t))); + CUDA_SAFE_CALL(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t))); init[thr_id] = true; } - if (opt_debug && throughput < (TPB * 4096)) - applog(LOG_DEBUG, "throughput=%u, start=%x, max=%x", throughput, first_nonce, max_nonce); - - blake256_cpu_setBlock_80(pdata, ptarget); + blake256_cpu_setBlock_80(pdata, ptarget, blakerounds); #if USE_CACHE crcsum = crc32_u32t(pdata, 64); #endif do { // GPU HASH - uint32_t foundNonce = blake256_cpu_hash_80(thr_id, throughput, pdata[19], blakerounds, crcsum); + uint32_t foundNonce = blake256_cpu_hash_80(thr_id, throughput, pdata[19], crcsum); if (foundNonce != MAXU) { uint32_t endiandata[20]; diff --git a/miner.h b/miner.h index d6127e2..7357978 100644 --- a/miner.h +++ b/miner.h @@ -243,7 +243,7 @@ extern int scanhash_anime(int thr_id, uint32_t *pdata, extern int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, - unsigned long *hashes_done, uint32_t blakerounds); + unsigned long *hashes_done, int8_t blakerounds); extern int scanhash_fresh(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, @@ -437,7 +437,7 @@ void applog_compare_hash(unsigned char *hash, unsigned char *hash2); void print_hash_tests(void); void animehash(void *state, const void *input); -void blake256hash(void *output, const void *input, int rounds); +void blake256hash(void *output, const void *input, int8_t rounds); void deephash(void *state, const void *input); void doomhash(void *state, const void *input); void fresh_hash(void *state, const void *input);