From 93bb428bdf9cd57e86177f32f0847bb262d649e6 Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Sun, 2 Nov 2014 16:23:01 +0100 Subject: [PATCH] blake: rewrite the cache system Unlike other hash algos, blake256 compute the hash with blocks of 64 bytes. We can do the first part on the cpu, only the 4 last int32 are computed on gpu (including the tested nonce) Previous method was also using this kind of cache with a crc. Blake Hash Speed: +5% --- blake32.cu | 219 ++++++++++++++++++++++++++++++++++++++-------- configure.ac | 2 +- cpu-miner.c | 3 +- cpuminer-config.h | 6 +- util.c | 24 +++++ 5 files changed, 213 insertions(+), 41 deletions(-) diff --git a/blake32.cu b/blake32.cu index 025a2af..be43604 100644 --- a/blake32.cu +++ b/blake32.cu @@ -1,11 +1,15 @@ /** * Blake-256 Cuda Kernel (Tested on SM 5.0) * - * Tanguy Pruvot - Aug. 2014 + * Tanguy Pruvot - Nov. 2014 */ +#define PRECALC64 1 + #include "miner.h" +typedef unsigned char uchar; + extern "C" { #include "sph/sph_blake.h" #include @@ -15,9 +19,7 @@ extern "C" { /* threads per block */ #define TPB 128 -/* crc32.c */ -extern "C" uint32_t crc32_u32t(const uint32_t *buf, size_t size); - +/* added in sph_blake.c */ extern "C" int blake256_rounds = 14; /* hash by cpu with blake 256 */ @@ -26,11 +28,11 @@ extern "C" void blake256hash(void *output, const void *input, int8_t rounds = 14 unsigned char hash[64]; sph_blake256_context ctx; - /* in sph_blake.c */ blake256_rounds = rounds; sph_blake256_init(&ctx); sph_blake256(&ctx, input, 80); + sph_blake256_close(&ctx, hash); memcpy(output, hash, 32); } @@ -41,11 +43,20 @@ extern "C" void blake256hash(void *output, const void *input, int8_t rounds = 14 // in cpu-miner.c extern bool opt_n_threads; -extern bool opt_benchmark; +extern bool opt_tracegpu; extern int device_map[8]; __constant__ +#if PRECALC64 +static uint32_t __align__(32) c_data[11]; +#else static uint32_t __align__(32) c_data[20]; +/* midstate hash cache, this algo is run on 2 parts */ +__device__ static uint32_t cache[8]; +__device__ static uint32_t prevsum = 0; +/* crc32.c */ +extern "C" uint32_t crc32_u32t(const uint32_t *buf, size_t size); +#endif /* 8 adapters max (-t threads) */ static uint32_t *d_resNonce[8]; @@ -53,14 +64,7 @@ static uint32_t *h_resNonce[8]; /* max count of found nounces in one call */ #define NBN 2 -static uint32_t extra_results[NBN-1] = { MAXU }; - -#define USE_CACHE 1 -/* midstate hash cache, this algo is run on 2 parts */ -#if USE_CACHE -__device__ static uint32_t cache[8]; -__device__ static uint32_t prevsum = 0; -#endif +static uint32_t extra_results[NBN] = { MAXU }; /* prefer uint32_t to prevent size conversions = speed +5/10 % */ __constant__ @@ -84,6 +88,7 @@ const uint32_t host_sigma[16][16] = { { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 } }; +#if !PRECALC64 __device__ __constant__ static const uint32_t __align__(32) c_IV256[8] = { SPH_C32(0x6A09E667), SPH_C32(0xBB67AE85), @@ -91,6 +96,7 @@ static const uint32_t __align__(32) c_IV256[8] = { SPH_C32(0x510E527F), SPH_C32(0x9B05688C), SPH_C32(0x1F83D9AB), SPH_C32(0x5BE0CD19) }; +#endif __device__ __constant__ static const uint32_t __align__(32) c_u256[16] = { @@ -105,8 +111,8 @@ static const uint32_t __align__(32) c_u256[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]; \ + const uint32_t idx1 = c_sigma[r][x]; \ + const uint32_t idx2 = c_sigma[r][x+1]; \ v[a] += (m[idx1] ^ c_u256[idx2]) + v[b]; \ v[d] = SPH_ROTL32(v[d] ^ v[a], 16); \ v[c] += v[d]; \ @@ -139,7 +145,11 @@ void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, co m[3] = block[3]; for (uint32_t i = 4; i < 16; i++) { +#if PRECALC64 + m[i] = c_Padding[i]; +#else m[i] = (T0 == 0x200) ? block[i] : c_Padding[i]; +#endif } //#pragma unroll 8 @@ -156,7 +166,7 @@ void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, co v[14] = c_u256[6]; v[15] = c_u256[7]; - for (int i = 0; i < rounds; i++) { + for (int r = 0; r < rounds; r++) { /* column step */ GS(0, 4, 0x8, 0xC, 0x0); GS(1, 5, 0x9, 0xD, 0x2); @@ -168,14 +178,20 @@ void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, co GS(2, 7, 0x8, 0xD, 0xC); GS(3, 4, 0x9, 0xE, 0xE); } - +#if PRECALC64 + // only compute h6 & 7 + h[6U] ^= v[6U] ^ v[14U]; + h[7U] ^= v[7U] ^ v[15U]; +#else //#pragma unroll 16 for (uint32_t i = 0; i < 16; i++) { uint32_t j = i % 8U; h[j] ^= v[i]; } +#endif } +#if !PRECALC64 /* original method */ __global__ void blake256_gpu_hash_80(const uint32_t threads, const uint32_t startNonce, uint32_t *resNounce, const uint64_t highTarget, const int crcsum, const int rounds) @@ -191,9 +207,6 @@ void blake256_gpu_hash_80(const uint32_t threads, const uint32_t startNonce, uin h[i] = c_IV256[i]; } -#if !USE_CACHE - blake256_compress(h, c_data, 512); -#else if (crcsum != prevsum) { prevsum = crcsum; blake256_compress(h, c_data, 512, rounds); @@ -207,8 +220,8 @@ void blake256_gpu_hash_80(const uint32_t threads, const uint32_t startNonce, uin h[i] = cache[i]; } } -#endif - // ------ Close: Bytes 64 to 80 ------ + + // ------ Close: Bytes 64 to 80 ------ uint32_t ending[4]; ending[0] = c_data[16]; @@ -273,15 +286,131 @@ void blake256_cpu_setBlock_80(uint32_t *pdata, const uint32_t *ptarget) CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_data, data, sizeof(data), 0, cudaMemcpyHostToDevice)); CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_sigma, host_sigma, sizeof(host_sigma), 0, cudaMemcpyHostToDevice)); } +#else + +/* ############################################################################################################################### */ +/* Precalculated 1st 64-bytes block (midstate) method */ + +__global__ +void blake256_gpu_hash_16(const uint32_t threads, const uint32_t startNonce, uint32_t *resNounce, + const uint64_t highTarget, const int rounds, const bool trace) +{ + uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); + if (thread < threads) + { + const uint32_t nounce = startNonce + thread; + uint32_t h[8]; + + #pragma unroll + for(int i=0; i<8; i++) { + h[i] = c_data[i]; + } + + // ------ Close: Bytes 64 to 80 ------ + + uint32_t ending[4]; + ending[0] = c_data[8]; + ending[1] = c_data[9]; + ending[2] = c_data[10]; + ending[3] = nounce; /* our tested value */ + + blake256_compress(h, ending, 640, rounds); + + //if (h[7] == 0 && high64 <= highTarget) { + if (h[7] == 0) { +#if NBN == 2 + /* keep the smallest nounce, + extra one if found */ + if (resNounce[0] > nounce) { + // printf("%llx %llx \n", high64, highTarget); + resNounce[1] = resNounce[0]; + resNounce[0] = nounce; + } + else + resNounce[1] = nounce; +#else + resNounce[0] = nounce; +#endif + if (trace) { +#ifdef _DEBUG + printf("tgt: %16llx\n", highTarget); + uint64_t high64 = ((uint64_t*)h)[3]; + printf("gpu: %16llx\n", high64); + printf("gpu: %16llx\n", cuda_swab64(h64)); + printf("gpu: %16x\n", cuda_swab32(h[6])); + printf("gpu: %08x %08x\n", h[6], h[7]); +#endif + } + } + } +} + +__host__ +static uint32_t blake256_cpu_hash_16(const int thr_id, const uint32_t threads, const uint32_t startNonce, const uint64_t highTarget, + const int8_t rounds) +{ + const int threadsperblock = TPB; + uint32_t result = MAXU; + + dim3 grid((threads + threadsperblock-1)/threadsperblock); + dim3 block(threadsperblock); + + /* Check error on Ctrl+C or kill to prevent segfaults on exit */ + if (cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t)) != cudaSuccess) + return result; + + blake256_gpu_hash_16 <<>> (threads, startNonce, d_resNonce[thr_id], highTarget, (int) rounds, opt_tracegpu); + cudaDeviceSynchronize(); + if (cudaSuccess == cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) { + //cudaThreadSynchronize(); /* seems no more required */ + result = h_resNonce[thr_id][0]; + for (int n=0; n < (NBN-1); n++) + extra_results[n] = h_resNonce[thr_id][n+1]; + } + return result; +} + +__host__ +static void blake256mid(uint32_t *output, const uint32_t *input, int8_t rounds = 14) +{ + sph_blake256_context ctx; + + /* in sph_blake.c */ + blake256_rounds = rounds; + + sph_blake256_init(&ctx); + sph_blake256(&ctx, input, 64); + + memcpy(output, (uchar*)ctx.H, 32); +} + +__host__ +void blake256_cpu_setBlock_16(uint32_t *penddata, const uint32_t *midstate, const uint32_t *ptarget) +{ + uint32_t data[11]; + memcpy(data, midstate, 32); + data[8] = penddata[0]; + data[9] = penddata[1]; + data[10]= penddata[2]; + CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_data, data, 32 + 12, 0, cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_sigma, host_sigma, sizeof(host_sigma), 0, cudaMemcpyHostToDevice)); +} +#endif extern "C" int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *ptarget, 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 }; - uint32_t throughput = min(TPB * 4096, max_nonce - first_nonce); - uint64_t targetHigh = ((uint64_t*)ptarget)[3]; - uint32_t crcsum = MAXU; + uint64_t targetHigh = ((uint64_t*)ptarget)[3]; // 0x00000000.0fffffff + uint32_t endiandata[20]; +#if PRECALC64 + uint32_t midstate[8]; +#else + uint32_t crcsum; +#endif + /* todo: -i param */ + uint32_t throughput = min(256 * 4096, max_nonce - first_nonce); + int rc = 0; #if NBN > 1 @@ -298,7 +427,14 @@ extern "C" int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *pt #endif if (opt_benchmark) - ((uint32_t*)ptarget)[7] = 0x00000f; + targetHigh = 0x1ULL << 32; + + if (opt_tracegpu) { + /* test call from util.c */ + throughput = 1; + for (int k = 0; k < 20; k++) + pdata[k] = swab32(pdata[k]); + } if (!init[thr_id]) { if (opt_n_threads > 1) { @@ -309,28 +445,38 @@ extern "C" int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *pt init[thr_id] = true; } +#if PRECALC64 + for (int k = 0; k < 16; k++) + be32enc(&endiandata[k], pdata[k]); + blake256mid(midstate, endiandata, blakerounds); + blake256_cpu_setBlock_16(&pdata[16], midstate, ptarget); +#else blake256_cpu_setBlock_80(pdata, ptarget); -#if USE_CACHE crcsum = crc32_u32t(pdata, 64); -#endif +#endif /* PRECALC64 */ do { - // GPU HASH - uint32_t foundNonce = blake256_cpu_hash_80(thr_id, throughput, pdata[19], targetHigh, crcsum, blakerounds); + uint32_t foundNonce = +#if PRECALC64 + // GPU HASH (second block only, first is midstate) + blake256_cpu_hash_16(thr_id, throughput, pdata[19], targetHigh, blakerounds); +#else + // GPU FULL HASH + blake256_cpu_hash_80(thr_id, throughput, pdata[19], targetHigh, crcsum, blakerounds); +#endif if (foundNonce != MAXU) { - uint32_t endiandata[20]; uint32_t vhashcpu[8]; - uint32_t Htarg = ptarget[6]; + uint32_t Htarg = (uint32_t)targetHigh; for (int k=0; k < 19; k++) be32enc(&endiandata[k], pdata[k]); be32enc(&endiandata[19], foundNonce); - blake256hash(vhashcpu, endiandata, blakerounds); - if (vhashcpu[6] <= Htarg || cuda_swab32(vhashcpu[6]) <= Htarg /*&& fulltest(vhashcpu, ptarget)*/) + //applog(LOG_BLUE, "%08x %16llx", vhashcpu[6], targetHigh); + if (vhashcpu[6] <= Htarg /*&& fulltest(vhashcpu, ptarget)*/) { pdata[19] = foundNonce; rc = 1; @@ -338,6 +484,7 @@ extern "C" int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *pt if (extra_results[0] != MAXU) { // Rare but possible if the throughput is big be32enc(&endiandata[19], extra_results[0]); + blake256hash(vhashcpu, endiandata, blakerounds); if (vhashcpu[6] <= Htarg /* && fulltest(vhashcpu, ptarget) */) { applog(LOG_NOTICE, "GPU found more than one result " CL_GRN "yippee!"); @@ -368,6 +515,6 @@ extern "C" int scanhash_blake256(int thr_id, uint32_t *pdata, const uint32_t *pt } while (!work_restart[thr_id].restart); exit_scan: - *hashes_done = pdata[19] - first_nonce + 1; + *hashes_done = pdata[19] - first_nonce + 1; // (+1 to prevent locks) return rc; } diff --git a/configure.ac b/configure.ac index a16a6e3..ea8032b 100644 --- a/configure.ac +++ b/configure.ac @@ -1,4 +1,4 @@ -AC_INIT([ccminer], [1.4.6]) +AC_INIT([ccminer], [1.4.7]) AC_PREREQ([2.59c]) AC_CANONICAL_SYSTEM diff --git a/cpu-miner.c b/cpu-miner.c index 6b9ed7b..e352b98 100644 --- a/cpu-miner.c +++ b/cpu-miner.c @@ -188,6 +188,7 @@ static const char *algo_names[] = { }; bool opt_debug = false; +bool opt_tracegpu = false; bool opt_protocol = false; bool opt_benchmark = false; bool want_longpoll = true; @@ -1534,7 +1535,7 @@ out: return NULL; } -#define PROGRAM_VERSION "1.4.6" +#define PROGRAM_VERSION "1.4.7" static void show_version_and_exit(void) { printf("%s v%s\n" diff --git a/cpuminer-config.h b/cpuminer-config.h index 53b1743..310b265 100644 --- a/cpuminer-config.h +++ b/cpuminer-config.h @@ -156,7 +156,7 @@ #define PACKAGE_NAME "ccminer" /* Define to the full name and version of this package. */ -#define PACKAGE_STRING "ccminer 1.4.6" +#define PACKAGE_STRING "ccminer 1.4.7" /* Define to the one symbol short name of this package. */ #define PACKAGE_TARNAME "ccminer" @@ -165,7 +165,7 @@ #define PACKAGE_URL "" /* Define to the version of this package. */ -#define PACKAGE_VERSION "1.4.6" +#define PACKAGE_VERSION "1.4.7" /* If using the C implementation of alloca, define if you know the direction of stack growth for your system; otherwise it will be @@ -188,7 +188,7 @@ #define USE_XOP 1 /* Version number of package */ -#define VERSION "1.4.6" +#define VERSION "1.4.7" /* Define curl_free() as free() if our version of curl lacks curl_free. */ /* #undef curl_free */ diff --git a/util.c b/util.c index eb973cb..41cbb10 100644 --- a/util.c +++ b/util.c @@ -1462,11 +1462,33 @@ extern void applog_hash(unsigned char *hash) #define printpfx(n,h) \ printf("%s%12s%s: %s\n", CL_BLU, n, CL_N, format_hash(s, h)) +extern bool opt_tracegpu; +void do_gpu_tests(void) +{ +#ifdef _DEBUG + unsigned long done; + char s[128] = { '\0' }; + unsigned char buf[128], hash[128]; + uint32_t tgt[8] = { 0 }; + memset(buf, 0, sizeof buf); + buf[0] = 1; buf[64] = 2; + opt_tracegpu = true; + work_restart = (struct work_restart*) malloc(sizeof(struct work_restart)); + work_restart[0].restart = 1; + tgt[6] = 0xffff; + scanhash_blake256(0, (uint32_t*)buf, tgt, 1, &done, 14); + free(work_restart); + work_restart = NULL; + opt_tracegpu = false; +#endif +} + void print_hash_tests(void) { char s[128] = {'\0'}; unsigned char buf[128], hash[128]; memset(buf, 0, sizeof buf); + // buf[0] = 1; buf[64] = 2; printf(CL_WHT "CPU HASH ON EMPTY BUFFER RESULTS:" CL_N "\n"); @@ -1482,6 +1504,8 @@ void print_hash_tests(void) blake256hash(&hash[0], &buf[0], 14); printpfx("blake", hash); + do_gpu_tests(); + memset(hash, 0, sizeof hash); deephash(&hash[0], &buf[0]); printpfx("deep", hash);