From 3341e0324ff083112c9e4af8d7b7112a7243b9ee Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Thu, 4 Sep 2014 16:17:11 +0200 Subject: [PATCH] blake: speed +10%, no more size conversions --- blake32.cu | 91 +++++++++++++++++++++++++---------------------------- cpu-miner.c | 2 +- 2 files changed, 44 insertions(+), 49 deletions(-) diff --git a/blake32.cu b/blake32.cu index ccba68a..e0e6814 100644 --- a/blake32.cu +++ b/blake32.cu @@ -31,44 +31,45 @@ extern "C" void blake32hash(void *output, const void *input) // in cpu-miner.c extern bool opt_n_threads; extern bool opt_benchmark; -//extern bool opt_debug; extern int device_map[8]; extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); __constant__ -static uint32_t c_Target[8]; +static uint32_t __align__(32) c_PaddedMessage80[32]; // padded message (80 bytes + padding) __constant__ -static uint32_t __align__(32) c_PaddedMessage80[32]; // padded message (80 bytes + padding) +static uint32_t __align__(32) c_Target[8]; + +#define MAXU 0xffffffffU static uint32_t *d_resNounce[8]; static uint32_t *h_resNounce[8]; __constant__ -static uint8_t c_sigma[16][16]; -const uint8_t host_sigma[16][16] = -{ - { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }, - {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 }, - {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 }, - { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 }, - { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 }, - { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 }, - {12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 }, - {13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 }, - { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 }, - {10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13 , 0 }, - { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }, - {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 }, - {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 }, - { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 }, - { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 }, - { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 } +static uint32_t __align__(32) c_sigma[16][16]; +/* prefer uint32_t to prevent size conversions = speed +5/10 % */ +const uint32_t host_sigma[16][16] = { + { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }, + {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 }, + {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 }, + { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 }, + { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 }, + { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 }, + {12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 }, + {13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 }, + { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 }, + {10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13 , 0 }, + { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }, + {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 }, + {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 }, + { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 }, + { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 }, + { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 } }; __device__ __constant__ -static const uint32_t c_IV256[8] = { +static const uint32_t __align__(32) c_IV256[8] = { SPH_C32(0x6A09E667), SPH_C32(0xBB67AE85), SPH_C32(0x3C6EF372), SPH_C32(0xA54FF53A), SPH_C32(0x510E527F), SPH_C32(0x9B05688C), @@ -76,8 +77,7 @@ static const uint32_t c_IV256[8] = { }; __device__ __constant__ - -static const uint32_t c_u256[16] = { +static const uint32_t __align__(32) c_u256[16] = { SPH_C32(0x243F6A88), SPH_C32(0x85A308D3), SPH_C32(0x13198A2E), SPH_C32(0x03707344), SPH_C32(0xA4093822), SPH_C32(0x299F31D0), @@ -112,13 +112,15 @@ static const uint32_t c_u256[16] = { } while (0) #endif -#define GS(a,b,c,d,e) { \ - v[a] += (m[sigma[i][e]] ^ u256[sigma[i][e+1]]) + v[b]; \ - v[d] = SPH_ROTR32(v[d] ^ v[a], 16); \ +#define GS(a,b,c,d,x) { \ + const uint32_t idx1 = c_sigma[i][x]; \ + const uint32_t idx2 = c_sigma[i][x+1]; \ + v[a] += (m[idx1] ^ u256[idx2]) + v[b]; \ + v[d] = SPH_ROTL32(v[d] ^ v[a], 16); \ v[c] += v[d]; \ v[b] = SPH_ROTR32(v[b] ^ v[c], 12); \ \ - v[a] += (m[sigma[i][e+1]] ^ u256[sigma[i][e]]) + v[b]; \ + v[a] += (m[idx2] ^ u256[idx1]) + v[b]; \ v[d] = SPH_ROTR32(v[d] ^ v[a], 8); \ v[c] += v[d]; \ v[b] = SPH_ROTR32(v[b] ^ v[c], 7); \ @@ -127,11 +129,13 @@ static const uint32_t c_u256[16] = { #define BLAKE256_ROUNDS 14 __device__ static -void blake256_compress(uint32_t *h, uint32_t *block, uint8_t ((*sigma)[16]), const uint32_t *u256, const uint32_t T0, uint8_t nullt = 1) +void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0) { uint32_t /* __align__(8) */ v[16]; uint32_t /* __align__(8) */ m[16]; + const uint32_t* u256 = c_u256; + //#pragma unroll for (int i = 0; i < 16; ++i) { m[i] = block[i]; @@ -170,16 +174,6 @@ void blake256_compress(uint32_t *h, uint32_t *block, uint8_t ((*sigma)[16]), con h[i % 8] ^= v[i]; } -#if __CUDA_ARCH__ >= 200 -/* memory should be aligned to use __nvvm_memset */ -#if (__NV_POINTER_SIZE == 64) -# define SZCT uint64_t -#else -# define SZCT uint32_t -#endif -extern __device__ __device_builtin__ void __nvvm_memset(uint8_t *, unsigned char, SZCT, int); -#endif - __global__ void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resNounce) { @@ -194,7 +188,7 @@ void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resN for(int i=0; i<8; i++) h[i] = c_IV256[i]; - blake256_compress(h, c_PaddedMessage80, c_sigma, c_u256, 0x200); /* 512 = 0x200 */ + blake256_compress(h, c_PaddedMessage80, 0x200); /* 512 = 0x200 */ // ------ Close: Bytes 64 to 80 ------ @@ -217,7 +211,7 @@ void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resN msg[14] = 0; msg[15] = 0x280; - blake256_compress(h, msg, c_sigma, c_u256, 0x280); + blake256_compress(h, msg, 0x280); for (int i = 7; i >= 0; i--) { uint32_t hash = cuda_swab32(h[i]); @@ -239,17 +233,18 @@ __host__ uint32_t blake256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce) { const int threadsperblock = TPB; + uint32_t result = MAXU; dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 block(threadsperblock); size_t shared_size = 0; - uint32_t result = 0xffffffffU; - cudaMemset(d_resNounce[thr_id], 0xff, sizeof(uint32_t)); + /* Check error on Ctrl+C or kill to prevent segfaults on exit */ + if (cudaMemset(d_resNounce[thr_id], 0xff, sizeof(uint32_t)) != cudaSuccess) + return result; blake256_gpu_hash_80<<>>(threads, startNounce, d_resNounce[thr_id]); - MyStreamSynchronize(NULL, 1, thr_id); - + cudaDeviceSynchronize(); if (cudaSuccess == cudaMemcpy(h_resNounce[thr_id], d_resNounce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost)) { cudaThreadSynchronize(); result = *h_resNounce[thr_id]; @@ -258,7 +253,7 @@ uint32_t blake256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce } __host__ -void blake256_cpu_setBlock_80(uint32_t *pdata, const void *ptarget) +void blake256_cpu_setBlock_80(uint32_t *pdata, const uint32_t *ptarget) { uint32_t PaddedMessage[32]; memcpy(PaddedMessage, pdata, 80); @@ -291,7 +286,7 @@ extern "C" int scanhash_blake32(int thr_id, uint32_t *pdata, const uint32_t *pta if (throughput < (TPB * 2048)) applog(LOG_WARNING, "throughput=%u, start=%x, max=%x", throughput, first_nonce, max_nonce); - blake256_cpu_setBlock_80(pdata, (void*)ptarget); + blake256_cpu_setBlock_80(pdata, ptarget); do { // GPU HASH diff --git a/cpu-miner.c b/cpu-miner.c index d0dc418..dc86a3e 100644 --- a/cpu-miner.c +++ b/cpu-miner.c @@ -962,7 +962,7 @@ static void *miner_thread(void *userdata) max64 = 0x1fffLL; break; case ALGO_BLAKE: - /* based on the 750Ti hashrate */ + /* based on the 750Ti hashrate (100kH) */ max64 = 0x3ffffffLL; break; default: