From 5a7eea9c94c51d9cc6b023b77d515ffdd7bc9202 Mon Sep 17 00:00:00 2001 From: KlausT Date: Sun, 1 Feb 2015 21:05:48 +0100 Subject: [PATCH 01/23] update help text --- README.txt | 1 + ccminer.cpp | 1 + 2 files changed, 2 insertions(+) diff --git a/README.txt b/README.txt index 910aad4850..d66f80a1fd 100644 --- a/README.txt +++ b/README.txt @@ -62,6 +62,7 @@ its command line interface and options. -a, --algo=ALGO specify the algorithm to use anime use to mine Animecoin + bitcoin use to mine Bitcoin blake use to mine NEOS (Blake 256) blakecoin use to mine Old Blake 256 deep use to mine Deepcoin diff --git a/ccminer.cpp b/ccminer.cpp index 7bd3722b1a..fdfee79b75 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -222,6 +222,7 @@ Usage: " PROGRAM_NAME " [OPTIONS]\n\ Options:\n\ -a, --algo=ALGO specify the hash algorithm to use\n\ anime Animecoin\n\ + bitcoin Bitcoin\n\ blake Blake 256 (SFR/NEOS)\n\ blakecoin Fast Blake 256 (8 rounds)\n\ deep Deepcoin\n\ From 4af01496eac392514cd830a05ca0898afe7dc5cd Mon Sep 17 00:00:00 2001 From: KlausT Date: Sun, 1 Feb 2015 23:52:14 +0100 Subject: [PATCH 02/23] bitcoin: fix wrong endianness --- bitcoin.cu | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/bitcoin.cu b/bitcoin.cu index 6f8b2b1107..77ef338d47 100644 --- a/bitcoin.cu +++ b/bitcoin.cu @@ -112,14 +112,14 @@ void bitcoin_hash(uint32_t *output, const uint32_t *data, uint32_t nonce, const b = a; a = t1 + t2; } - output[0] = a + hc[0]; - output[1] = b + hc[1]; - output[2] = c + hc[2]; - output[3] = d + hc[3]; - output[4] = e + hc[4]; - output[5] = f + hc[5]; - output[6] = g + hc[6]; - output[7] = h + hc[7]; + be32enc(&output[0], a + hc[0]); + be32enc(&output[1], b + hc[1]); + be32enc(&output[2], c + hc[2]); + be32enc(&output[3], d + hc[3]); + be32enc(&output[4], e + hc[4]); + be32enc(&output[5], f + hc[5]); + be32enc(&output[6], g + hc[6]); + be32enc(&output[7], h + hc[7]); } static bool init[MAX_GPUS] = { 0 }; From 39810fbea48bde55f8338e156cde171b511b362a Mon Sep 17 00:00:00 2001 From: KlausT Date: Tue, 3 Feb 2015 05:37:38 +0100 Subject: [PATCH 03/23] optimization for Groestlcoin and Diamond-Groestl --- bitslice_transformations_quad.cu | 32 ++++++++++++++ cuda_groestlcoin.cu | 73 ++++++++++++-------------------- groestlcoin.cpp | 49 +++++++++++++++------ 3 files changed, 93 insertions(+), 61 deletions(-) diff --git a/bitslice_transformations_quad.cu b/bitslice_transformations_quad.cu index ddbeb1aa81..57b05528e6 100644 --- a/bitslice_transformations_quad.cu +++ b/bitslice_transformations_quad.cu @@ -166,3 +166,35 @@ void from_bitslice_quad(const uint32_t *const __restrict__ input, uint32_t *cons } */ } +__device__ __forceinline__ +void from_bitslice_quad_final(const uint32_t *const __restrict__ input, uint32_t *const __restrict__ output) +{ + + uint32_t t; + + output[0] = __byte_perm(input[0], input[4], 0x7531); + output[2] = __byte_perm(input[1], input[5], 0x7531); + output[8] = __byte_perm(input[2], input[6], 0x7531); + output[10] = __byte_perm(input[3], input[7], 0x7531); + + SWAP1(output[0], output[2]); + SWAP1(output[8], output[10]); + + SWAP2(output[2], output[10]); + + output[6] = __byte_perm(output[2], output[10], 0x5410); + output[10] = __byte_perm(output[2], output[10], 0x7632); + output[2] = output[6]; + + SWAP4(output[2], output[10]); + + output[6] = output[2]; + + if (threadIdx.x & 1) + { + output[6] = __byte_perm(output[2], 0, 0x3232); + } + + output[6] = __byte_perm(output[6], __shfl((int)output[6], (threadIdx.x + 1) & 3, 4), 0x7632); + output[6 + 1] = __shfl((int)output[6], (threadIdx.x + 2) & 3, 4); +} diff --git a/cuda_groestlcoin.cu b/cuda_groestlcoin.cu index ce8ea18a5b..3e0d1bf6f9 100644 --- a/cuda_groestlcoin.cu +++ b/cuda_groestlcoin.cu @@ -22,7 +22,7 @@ __global__ __launch_bounds__(256, 4) void groestlcoin_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t *resNounce) { // durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen - uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) / 4; + uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) >> 2; if (thread < threads) { // GROESTL @@ -38,52 +38,31 @@ void groestlcoin_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t to_bitslice_quad(paddedInput, msgBitsliced); uint32_t state[8]; - for (int round=0; round<2; round++) - { - groestl512_progressMessage_quad(state, msgBitsliced); - - if (round < 1) - { - // Verkettung zweier Runden inclusive Padding. - msgBitsliced[ 0] = __byte_perm(state[ 0], 0x00800100, 0x4341 + ((threadIdx.x & 3)==3)*0x2000); - msgBitsliced[ 1] = __byte_perm(state[ 1], 0x00800100, 0x4341); - msgBitsliced[ 2] = __byte_perm(state[ 2], 0x00800100, 0x4341); - msgBitsliced[ 3] = __byte_perm(state[ 3], 0x00800100, 0x4341); - msgBitsliced[ 4] = __byte_perm(state[ 4], 0x00800100, 0x4341); - msgBitsliced[ 5] = __byte_perm(state[ 5], 0x00800100, 0x4341); - msgBitsliced[ 6] = __byte_perm(state[ 6], 0x00800100, 0x4341); - msgBitsliced[7] = __byte_perm(state[7], 0x00800100, 0x4341 + ((threadIdx.x & 3) == 0) * 0x0010); - } - } - - // Nur der erste von jeweils 4 Threads bekommt das Ergebns-Hash - uint32_t out_state[16]; - from_bitslice_quad(state, out_state); + groestl512_progressMessage_quad(state, msgBitsliced); + + msgBitsliced[ 0] = __byte_perm(state[ 0], 0x00800100, 0x4341 + ((threadIdx.x & 3)==3)*0x2000); + msgBitsliced[ 1] = __byte_perm(state[ 1], 0x00800100, 0x4341); + msgBitsliced[ 2] = __byte_perm(state[ 2], 0x00800100, 0x4341); + msgBitsliced[ 3] = __byte_perm(state[ 3], 0x00800100, 0x4341); + msgBitsliced[ 4] = __byte_perm(state[ 4], 0x00800100, 0x4341); + msgBitsliced[ 5] = __byte_perm(state[ 5], 0x00800100, 0x4341); + msgBitsliced[ 6] = __byte_perm(state[ 6], 0x00800100, 0x4341); + msgBitsliced[7] = __byte_perm(state[7], 0x00800100, 0x4341 + ((threadIdx.x & 3) == 0) * 0x0010); + + groestl512_progressMessage_quad(state, msgBitsliced); + + uint32_t out_state[16]; + from_bitslice_quad_final(state, out_state); if ((threadIdx.x & 3) == 0) { - int i, position = -1; - bool rc = true; - - #pragma unroll 8 - for (i = 7; i >= 0; i--) { - if (out_state[i] > pTarget[i]) { - if(position < i) { - position = i; - rc = false; - } - } - if (out_state[i] < pTarget[i]) { - if(position < i) { - position = i; - rc = true; - } - } - } - - if(rc == true) - if(resNounce[0] > nounce) - resNounce[0] = nounce; + if (out_state[7] <= pTarget[7]) + { + uint32_t tmp = atomicExch(resNounce, nounce); + if (tmp != 0xffffffff) + resNounce[1] = tmp; + } + } } } @@ -94,7 +73,7 @@ __host__ void groestlcoin_cpu_init(int thr_id, uint32_t threads) CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); // Speicher für Gewinner-Nonce belegen - cudaMalloc(&d_resultNonce[thr_id], sizeof(uint32_t)); + cudaMalloc(&d_resultNonce[thr_id], 2 * sizeof(uint32_t)); } __host__ void groestlcoin_cpu_setBlock(int thr_id, void *data, void *pTargetIn) @@ -136,8 +115,8 @@ __host__ void groestlcoin_cpu_hash(int thr_id, uint32_t threads, uint32_t startN dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock)); dim3 block(threadsperblock); - cudaMemset(d_resultNonce[thr_id], 0xFF, sizeof(uint32_t)); + cudaMemset(d_resultNonce[thr_id], 0xFF, 2 * sizeof(uint32_t)); groestlcoin_gpu_hash_quad<<>>(threads, startNounce, d_resultNonce[thr_id]); - cudaMemcpy(nounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost); + cudaMemcpy(nounce, d_resultNonce[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost); } diff --git a/groestlcoin.cpp b/groestlcoin.cpp index f732be93c2..1a133cbce3 100644 --- a/groestlcoin.cpp +++ b/groestlcoin.cpp @@ -62,27 +62,48 @@ extern "C" int scanhash_groestlcoin(int thr_id, uint32_t *pdata, const uint32_t do { // GPU - uint32_t foundNounce = 0xFFFFFFFF; + uint32_t foundNounce[2]; const uint32_t Htarg = ptarget[7]; - groestlcoin_cpu_hash(thr_id, throughput, pdata[19], outputHash, &foundNounce); + groestlcoin_cpu_hash(thr_id, throughput, pdata[19], outputHash, foundNounce); - if(foundNounce < 0xffffffff) + if(foundNounce[0] < 0xffffffff) { uint32_t tmpHash[8]; - endiandata[19] = SWAP32(foundNounce); + endiandata[19] = SWAP32(foundNounce[0]); groestlhash(tmpHash, endiandata); - if (tmpHash[7] <= Htarg && fulltest(tmpHash, ptarget)) { - pdata[19] = foundNounce; - *hashes_done = foundNounce - start_nonce + 1; - free(outputHash); - return true; - } else { - applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU!", thr_id, foundNounce); - } - - foundNounce = 0xffffffff; + if (tmpHash[7] <= Htarg && fulltest(tmpHash, ptarget)) + { + int res = 1; + // check if there was some other ones... + *hashes_done = pdata[19] - start_nonce + throughput; + if (foundNounce[1] != 0xffffffff) + { + endiandata[19] = SWAP32(foundNounce[1]); + groestlhash(tmpHash, endiandata); + + if (tmpHash[7] <= Htarg && fulltest(tmpHash, ptarget)) + { + + pdata[21] = foundNounce[1]; + res++; + if (opt_benchmark) + applog(LOG_INFO, "GPU #%d Found second nounce %08x", thr_id, foundNounce[1]); + } + } + pdata[19] = foundNounce[0]; + if (opt_benchmark) + applog(LOG_INFO, "GPU #%d Found nounce %08x", thr_id, foundNounce[0]); + return res; + } + else + { + if (tmpHash[7] != Htarg) + { + applog(LOG_INFO, "GPU #%d: result for %08x does not validate on CPU!", thr_id, foundNounce[0]); + } + } } pdata[19] += throughput; From caea1ecf5a32f5f01b4754bfa6c401379fa9bf2e Mon Sep 17 00:00:00 2001 From: KlausT Date: Tue, 3 Feb 2015 08:07:45 +0100 Subject: [PATCH 04/23] check if cudaSetDevice returns an error --- Algo256/keccak256.cu | 2 +- bitcoin.cu | 2 +- cuda_groestlcoin.cu | 3 +++ cuda_myriadgroestl.cu | 2 +- cuda_nist5.cu | 2 +- lyra2/lyra2RE.cu | 2 +- quark/animecoin.cu | 2 +- quark/quarkcoin.cu | 2 +- qubit/deep.cu | 2 +- qubit/doom.cu | 2 +- qubit/qubit.cu | 2 +- x11/fresh.cu | 2 +- x11/s3.cu | 2 +- x11/x11.cu | 2 +- x13/x13.cu | 2 +- x15/x14.cu | 2 +- x15/x15.cu | 2 +- x17/x17.cu | 2 +- 18 files changed, 20 insertions(+), 17 deletions(-) diff --git a/Algo256/keccak256.cu b/Algo256/keccak256.cu index 5d4db89b46..2a84b0cb56 100644 --- a/Algo256/keccak256.cu +++ b/Algo256/keccak256.cu @@ -49,7 +49,7 @@ extern "C" int scanhash_keccak256(int thr_id, uint32_t *pdata, ((uint32_t*)ptarget)[7] = 0x0005; if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/bitcoin.cu b/bitcoin.cu index 77ef338d47..4a130232ab 100644 --- a/bitcoin.cu +++ b/bitcoin.cu @@ -137,7 +137,7 @@ int scanhash_bitcoin(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/cuda_groestlcoin.cu b/cuda_groestlcoin.cu index 3e0d1bf6f9..6baeb89b77 100644 --- a/cuda_groestlcoin.cu +++ b/cuda_groestlcoin.cu @@ -71,6 +71,9 @@ void groestlcoin_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t __host__ void groestlcoin_cpu_init(int thr_id, uint32_t threads) { CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); + cudaDeviceReset(); + cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); + cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); // Speicher für Gewinner-Nonce belegen cudaMalloc(&d_resultNonce[thr_id], 2 * sizeof(uint32_t)); diff --git a/cuda_myriadgroestl.cu b/cuda_myriadgroestl.cu index d404afe7b4..4ba7f9555c 100644 --- a/cuda_myriadgroestl.cu +++ b/cuda_myriadgroestl.cu @@ -269,7 +269,7 @@ __global__ void __launch_bounds__(256, 3) // Setup-Funktionen __host__ void myriadgroestl_cpu_init(int thr_id, uint32_t threads) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/cuda_nist5.cu b/cuda_nist5.cu index 437889b648..e61397971f 100644 --- a/cuda_nist5.cu +++ b/cuda_nist5.cu @@ -80,7 +80,7 @@ extern "C" int scanhash_nist5(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/lyra2/lyra2RE.cu b/lyra2/lyra2RE.cu index 083caeb734..8b167beca9 100644 --- a/lyra2/lyra2RE.cu +++ b/lyra2/lyra2RE.cu @@ -72,7 +72,7 @@ extern "C" int scanhash_lyra2(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/quark/animecoin.cu b/quark/animecoin.cu index 4994e0a831..96ab889068 100644 --- a/quark/animecoin.cu +++ b/quark/animecoin.cu @@ -175,7 +175,7 @@ extern "C" int scanhash_anime(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/quark/quarkcoin.cu b/quark/quarkcoin.cu index 18a67521a8..bbff75bd99 100644 --- a/quark/quarkcoin.cu +++ b/quark/quarkcoin.cu @@ -149,7 +149,7 @@ extern "C" int scanhash_quark(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/qubit/deep.cu b/qubit/deep.cu index bb0a2ad1fd..05dc92d488 100644 --- a/qubit/deep.cu +++ b/qubit/deep.cu @@ -67,7 +67,7 @@ extern "C" int scanhash_deep(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/qubit/doom.cu b/qubit/doom.cu index 93f46ef0d8..54a3a0e8b5 100644 --- a/qubit/doom.cu +++ b/qubit/doom.cu @@ -48,7 +48,7 @@ extern "C" int scanhash_doom(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/qubit/qubit.cu b/qubit/qubit.cu index 538fabe3c3..bfd80d7410 100644 --- a/qubit/qubit.cu +++ b/qubit/qubit.cu @@ -88,7 +88,7 @@ extern "C" int scanhash_qubit(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/x11/fresh.cu b/x11/fresh.cu index c641176743..decd1553a4 100644 --- a/x11/fresh.cu +++ b/x11/fresh.cu @@ -87,7 +87,7 @@ extern "C" int scanhash_fresh(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/x11/s3.cu b/x11/s3.cu index e4c7749f6b..a637f8711e 100644 --- a/x11/s3.cu +++ b/x11/s3.cu @@ -71,7 +71,7 @@ extern "C" int scanhash_s3(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/x11/x11.cu b/x11/x11.cu index 7d02f47107..b33b327ad4 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -153,7 +153,7 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/x13/x13.cu b/x13/x13.cu index 2953306de4..433c0ab4b2 100644 --- a/x13/x13.cu +++ b/x13/x13.cu @@ -166,7 +166,7 @@ extern "C" int scanhash_x13(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/x15/x14.cu b/x15/x14.cu index 447ecb9f33..ecb3318a09 100644 --- a/x15/x14.cu +++ b/x15/x14.cu @@ -164,7 +164,7 @@ extern "C" int scanhash_x14(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/x15/x15.cu b/x15/x15.cu index 811b7c51c6..80450afbda 100644 --- a/x15/x15.cu +++ b/x15/x15.cu @@ -176,7 +176,7 @@ extern "C" int scanhash_x15(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); diff --git a/x17/x17.cu b/x17/x17.cu index 97ab886ab6..9881c04ff5 100644 --- a/x17/x17.cu +++ b/x17/x17.cu @@ -192,7 +192,7 @@ extern "C" int scanhash_x17(int thr_id, uint32_t *pdata, if (!init[thr_id]) { - cudaSetDevice(device_map[thr_id]); + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); cudaDeviceReset(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); From 3f324de6ecb1e102b322e771198ff09909b2a24a Mon Sep 17 00:00:00 2001 From: KlausT Date: Tue, 3 Feb 2015 08:09:05 +0100 Subject: [PATCH 05/23] remove unnecessary cudaDeviceSynchronize --- Algo256/blake256.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/Algo256/blake256.cu b/Algo256/blake256.cu index 21b0276eae..8a3a9f1e23 100644 --- a/Algo256/blake256.cu +++ b/Algo256/blake256.cu @@ -339,7 +339,6 @@ static uint32_t blake256_cpu_hash_16(const int thr_id, const uint32_t threads, c 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)) { //cudaDeviceSynchronize(); /* seems no more required */ result = h_resNonce[thr_id][0]; From e559cdaa680851098116519a25b19db5dfb75bb4 Mon Sep 17 00:00:00 2001 From: KlausT Date: Tue, 3 Feb 2015 08:24:35 +0100 Subject: [PATCH 06/23] check for CUDA version 6.5 --- cuda.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cuda.cpp b/cuda.cpp index 95c8221bdf..ec31727fd4 100644 --- a/cuda.cpp +++ b/cuda.cpp @@ -32,7 +32,7 @@ int cuda_num_devices() } int maj = version / 1000, min = version % 100; // same as in deviceQuery sample - if (maj < 5 || (maj == 5 && min < 5)) + if (maj < 5 || (maj == 6 && min < 5)) { applog(LOG_ERR, "Driver does not support CUDA %d.%d API! Update your nVidia driver!", 5, 5); exit(1); From 084a6bec10d94d49c6527ac51219891c326ee0e1 Mon Sep 17 00:00:00 2001 From: KlausT Date: Tue, 3 Feb 2015 17:33:14 +0100 Subject: [PATCH 07/23] fix ROR2 for Kepler cards --- cuda_helper.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cuda_helper.h b/cuda_helper.h index b3834c24ec..41d312a5f3 100644 --- a/cuda_helper.h +++ b/cuda_helper.h @@ -518,8 +518,8 @@ __device__ __inline__ uint2 ROR2(const uint2 v, const int n) } else { - result.y = ((v.x >> (n - 32)) | (v.y << (32 - n))); - result.x = ((v.y >> (n - 32)) | (v.x << (32 - n))); + result.y = ((v.x >> (n - 32)) | (v.y << (64 - n))); + result.x = ((v.y >> (n - 32)) | (v.x << (64 - n))); } return result; } From 228a8c41718ff78fb12580e1cc577bb5570d2d2d Mon Sep 17 00:00:00 2001 From: KlausT Date: Tue, 3 Feb 2015 17:43:05 +0100 Subject: [PATCH 08/23] small change for uint2 shift --- cuda_helper.h | 32 +++++++++++++++----------------- 1 file changed, 15 insertions(+), 17 deletions(-) diff --git a/cuda_helper.h b/cuda_helper.h index 41d312a5f3..955b6c75f3 100644 --- a/cuda_helper.h +++ b/cuda_helper.h @@ -627,10 +627,10 @@ uint2 SWAPDWORDS2(uint2 value) return make_uint2(value.y, value.x); } -static __forceinline__ __device__ uint2 SHL2(uint2 a, int offset) +static __forceinline__ __device__ uint2 SHL2(const uint2 a, int offset) { -#if __CUDA_ARCH__ > 300 uint2 result; +#if __CUDA_ARCH__ > 300 if (offset<32) { asm("{\n\t" @@ -646,25 +646,24 @@ static __forceinline__ __device__ uint2 SHL2(uint2 a, int offset) "}\n\t" : "=r"(result.x), "=r"(result.y) : "r"(a.y), "r"(a.x), "r"(offset)); } - return result; #else if (offset<=32) { - a.y = (a.y << offset) | (a.x >> (32 - offset)); - a.x = (a.x << offset); + result.y = (a.y << offset) | (a.x >> (32 - offset)); + result.x = (a.x << offset); } else { - a.y = (a.x << (offset-32)); - a.x = 0; + result.y = (a.x << (offset - 32)); + result.x = 0; } - return a; #endif + return result; } -static __forceinline__ __device__ uint2 SHR2(uint2 a, int offset) +static __forceinline__ __device__ uint2 SHR2(const uint2 a, int offset) { - #if __CUDA_ARCH__ > 300 uint2 result; +#if __CUDA_ARCH__ > 300 if (offset<32) { asm("{\n\t" "shf.r.clamp.b32 %0,%2,%3,%4; \n\t" @@ -679,20 +678,19 @@ static __forceinline__ __device__ uint2 SHR2(uint2 a, int offset) "}\n\t" : "=r"(result.x), "=r"(result.y) : "r"(a.y), "r"(a.x), "r"(offset)); } - return result; #else if (offset<=32) { - a.x = (a.x >> offset) | (a.y << (32 - offset)); - a.y = (a.y >> offset); + result.x = (a.x >> offset) | (a.y << (32 - offset)); + result.y = (a.y >> offset); } else { - a.x = (a.y >> (offset - 32)); - a.y = 0; + result.x = (a.y >> (offset - 32)); + result.y = 0; } - return a; - #endif +#endif + return result; } static __device__ __forceinline__ uint64_t devectorizeswap(uint2 v) { return MAKE_ULONGLONG(cuda_swab32(v.y), cuda_swab32(v.x)); } From 18e39a2df63b20ca8c0643d6b895b2d4c05c30fc Mon Sep 17 00:00:00 2001 From: sp-hash Date: Tue, 3 Feb 2015 20:17:26 +0100 Subject: [PATCH 09/23] Bitcoin 17% faster on the 750ti --- bitcoin.cu | 2 +- cuda_bitcoin.cu | 10 ++++------ 2 files changed, 5 insertions(+), 7 deletions(-) diff --git a/bitcoin.cu b/bitcoin.cu index 4a130232ab..97b537bbcf 100644 --- a/bitcoin.cu +++ b/bitcoin.cu @@ -129,7 +129,7 @@ int scanhash_bitcoin(int thr_id, uint32_t *pdata, unsigned long *hashes_done) { const uint32_t first_nonce = pdata[19]; - uint32_t throughput = device_intensity(thr_id, __func__, 1U << 24); + uint32_t throughput = device_intensity(thr_id, __func__, 1U << 28); throughput = min(throughput, (max_nonce - first_nonce)); if (opt_benchmark) diff --git a/cuda_bitcoin.cu b/cuda_bitcoin.cu index 80c6dc6965..250392445a 100644 --- a/cuda_bitcoin.cu +++ b/cuda_bitcoin.cu @@ -13,15 +13,11 @@ __constant__ uint32_t pTarget[8]; static uint32_t *d_result[MAX_GPUS]; #define TPB 512 -#define NONCES_PER_THREAD 2048 +#define NONCES_PER_THREAD 32 -#if __CUDA_ARCH__ < 320 #define rrot(x, n) ((x >> n) | (x << (32 - n))) -#else -#define rrot(x, n) __funnelshift_r((x), (x), (n)) -#endif -__global__ __launch_bounds__(TPB, 1) +__global__ __launch_bounds__(TPB, 2) void bitcoin_gpu_hash(const uint32_t threads, const uint32_t startNounce, uint32_t *const result, const uint32_t t1c, const uint32_t t2c, const uint32_t w16, const uint32_t w16rot, const uint32_t w17, const uint32_t w17rot, const uint32_t b2, const uint32_t c2, const uint32_t d2, const uint32_t f2, const uint32_t g2, const uint32_t h2, const uint32_t ms0, const uint32_t ms1, const uint32_t ms2, const uint32_t ms3, const uint32_t ms4, const uint32_t ms5, const uint32_t ms6, const uint32_t ms7, const uint32_t compacttarget) { uint32_t threadindex = (blockDim.x * blockIdx.x + threadIdx.x); @@ -33,6 +29,8 @@ void bitcoin_gpu_hash(const uint32_t threads, const uint32_t startNounce, uint32 const uint32_t maxnonce = startNounce + threadindex + numberofthreads*NONCES_PER_THREAD - 1; const uint32_t threadindex = blockIdx.x*blockDim.x + threadIdx.x; + + #pragma unroll for (uint32_t nonce = startNounce + threadindex; nonce <= maxnonce; nonce += numberofthreads) { w[18] = (rrot(nonce, 7) ^ rrot(nonce, 18) ^ (nonce >> 3)) + w16rot; From ccce94e38c6cf301f8b51f5e303e4294a2e5c0c3 Mon Sep 17 00:00:00 2001 From: KlausT Date: Tue, 3 Feb 2015 21:26:52 +0100 Subject: [PATCH 10/23] change x11 default intensity --- x11/x11.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/x11/x11.cu b/x11/x11.cu index b33b327ad4..03d6561729 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -144,8 +144,9 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, const uint32_t first_nonce = pdata[19]; int intensity = 256 * 256 * 20; - if (device_sm[device_map[thr_id]] == 520) intensity = 256 * 256 * 41; + if (device_sm[device_map[thr_id]] == 520) intensity = 256 * 256 * 32; uint32_t throughput = device_intensity(thr_id, __func__, intensity); // 19=256*256*8; + throughput = min(throughput, (max_nonce - first_nonce)); if (opt_benchmark) From 883ef80389c29e9b12b64afc08508449f702047d Mon Sep 17 00:00:00 2001 From: KlausT Date: Thu, 5 Feb 2015 08:13:52 +0100 Subject: [PATCH 11/23] fix lyra2Re intensity bug --- lyra2/lyra2RE.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/lyra2/lyra2RE.cu b/lyra2/lyra2RE.cu index 8b167beca9..5d4c639b2f 100644 --- a/lyra2/lyra2RE.cu +++ b/lyra2/lyra2RE.cu @@ -64,7 +64,7 @@ extern "C" int scanhash_lyra2(int thr_id, uint32_t *pdata, { const uint32_t first_nonce = pdata[19]; unsigned int intensity = (device_sm[device_map[thr_id]] > 500) ? 256 * 256 * 25 : 256 * 256 * 14; - uint32_t throughput = device_intensity(thr_id, __func__, 1U << intensity); // 18=256*256*4; + uint32_t throughput = device_intensity(thr_id, __func__, intensity); // 18=256*256*4; throughput = min(throughput, (max_nonce - first_nonce)); if (opt_benchmark) From ff5eb40c5917cff78c59bc141e8b8be967f6fa5d Mon Sep 17 00:00:00 2001 From: sp-hash Date: Thu, 5 Feb 2015 00:13:34 +0100 Subject: [PATCH 12/23] Rolled back simd changes. Added better launch config. Simd +6.66% (750ti) --- x11/cuda_x11_simd512.cu | 8 ++--- x11/fresh.cu | 50 +++++------------------------- x11/x11.cu | 68 +++++++++-------------------------------- 3 files changed, 25 insertions(+), 101 deletions(-) diff --git a/x11/cuda_x11_simd512.cu b/x11/cuda_x11_simd512.cu index 17a24ff616..bf059a9818 100644 --- a/x11/cuda_x11_simd512.cu +++ b/x11/cuda_x11_simd512.cu @@ -5,7 +5,7 @@ // // STEP8_IF and STEP8_MAJ beinhalten je 2x 8-fach parallel Operations -#define TPB 64 +#define TPB 256 #include "cuda_helper.h" #include @@ -566,7 +566,7 @@ x11_simd512_gpu_expand_64(uint32_t threads, uint32_t startNounce, const uint64_t } } -__global__ void __launch_bounds__(TPB, 4) +__global__ void __launch_bounds__(TPB, 1) x11_simd512_gpu_compress1_64(uint32_t threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector, uint4 *g_fft4, uint32_t *g_state) { uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); @@ -580,7 +580,7 @@ x11_simd512_gpu_compress1_64(uint32_t threads, uint32_t startNounce, uint64_t *g Compression1(Hash, hashPosition, g_fft4, g_state); } } -__global__ void __launch_bounds__(TPB, 4) +__global__ void __launch_bounds__(TPB, 1) x11_simd512_gpu_compress2_64(uint32_t threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector, uint4 *g_fft4, uint32_t *g_state) { uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); @@ -595,7 +595,7 @@ x11_simd512_gpu_compress2_64(uint32_t threads, uint32_t startNounce, uint64_t *g } -__global__ void __launch_bounds__(TPB, 4) +__global__ void __launch_bounds__(TPB, 1) x11_simd512_gpu_compress_64_maxwell(uint32_t threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector, uint4 *g_fft4, uint32_t *g_state) { uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); diff --git a/x11/fresh.cu b/x11/fresh.cu index decd1553a4..b28a07b429 100644 --- a/x11/fresh.cu +++ b/x11/fresh.cu @@ -14,7 +14,6 @@ extern "C" { static uint32_t *d_hash[MAX_GPUS]; static uint32_t *h_found[MAX_GPUS]; -static uint32_t *h_found2[MAX_GPUS]; extern void x11_shavite512_setBlock_80(void *pdata); extern void x11_shavite512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash, int order); @@ -79,7 +78,7 @@ extern "C" int scanhash_fresh(int thr_id, uint32_t *pdata, const uint32_t first_nonce = pdata[19]; uint32_t endiandata[20]; - uint32_t throughput = device_intensity(thr_id, __func__, 1 << 20); + uint32_t throughput = device_intensity(thr_id, __func__, 1 << 19); throughput = min(throughput, (max_nonce - first_nonce)); if (opt_benchmark) @@ -92,12 +91,11 @@ extern "C" int scanhash_fresh(int thr_id, uint32_t *pdata, cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); - x11_simd512_cpu_init(thr_id, throughput>>1); + x11_simd512_cpu_init(thr_id, throughput); x11_echo512_cpu_init(thr_id, throughput); CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput + 4), 0); CUDA_CALL_OR_RET_X(cudaMallocHost(&(h_found[thr_id]), 4 * sizeof(uint32_t)), 0); - CUDA_CALL_OR_RET_X(cudaMallocHost(&(h_found2[thr_id]), 4 * sizeof(uint32_t)), 0); cuda_check_cpu_init(thr_id, throughput); @@ -117,14 +115,11 @@ extern "C" int scanhash_fresh(int thr_id, uint32_t *pdata, // GPU Hash - x11_shavite512_cpu_hash_80(thr_id, throughput >> 1, pdata[19] + (throughput >> 1), d_hash[thr_id] + (throughput * 8), order++); - x11_shavite512_cpu_hash_80(thr_id, throughput>>1, pdata[19], d_hash[thr_id], order++); - x11_simd512_cpu_hash_64(thr_id, throughput>>1, pdata[19], NULL, d_hash[thr_id], order++); - x11_simd512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput >> 1), NULL, d_hash[thr_id] + (throughput * 8), order++); - x11_shavite512_cpu_hash_64(thr_id, throughput>>1, pdata[19], NULL, d_hash[thr_id], order++); - x11_shavite512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput >> 1), NULL, d_hash[thr_id] + (throughput * 8), order++); - x11_simd512_cpu_hash_64(thr_id, throughput>>1, pdata[19], NULL, d_hash[thr_id], order++); - x11_echo512_cpu_hash_64_final(thr_id, throughput>>1, pdata[19], NULL, d_hash[thr_id], ptarget[7], h_found[thr_id], order++); + x11_shavite512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x11_echo512_cpu_hash_64_final(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], ptarget[7], h_found[thr_id], order++); cudaDeviceSynchronize(); if (h_found[thr_id][0] != 0xffffffff) @@ -153,37 +148,6 @@ extern "C" int scanhash_fresh(int thr_id, uint32_t *pdata, applog(LOG_INFO, "GPU #%d: result for %08x does not validate on CPU!", thr_id, h_found[thr_id][0]); } } - - x11_simd512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput >> 1), NULL, d_hash[thr_id] + (throughput * 8), order++); - x11_echo512_cpu_hash_64_final(thr_id, throughput >> 1, pdata[19] + (throughput >> 1), NULL, d_hash[thr_id] + (throughput * 8), ptarget[7], h_found2[thr_id], order++); - if (h_found2[thr_id][0] != 0xffffffff) - { - uint32_t vhash64[8]; - be32enc(&endiandata[19], h_found2[thr_id][0]); - fresh_hash(vhash64, endiandata); - - if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) - { - int res = 1; - *hashes_done = pdata[19] - first_nonce + throughput; - if (opt_benchmark) - applog(LOG_INFO, "GPU #%d Found nounce %08x", thr_id, h_found2[thr_id][0], vhash64[7], Htarg); - - if (h_found2[thr_id][1] != 0xffffffff) - { - pdata[21] = h_found2[thr_id][1]; - res++; - } - pdata[19] = h_found2[thr_id][0]; - return res; - } - else - { - applog(LOG_INFO, "GPU #%d: result for %08x does not validate on CPU!", thr_id, h_found2[thr_id][0]); - } - } - - pdata[19] += throughput; } while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > ((uint64_t)(pdata[19]) + (uint64_t)throughput))); diff --git a/x11/x11.cu b/x11/x11.cu index 03d6561729..fc0c0eeaf8 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -29,7 +29,6 @@ extern "C" uint32_t *d_hash[MAX_GPUS]; uint32_t *h_found[MAX_GPUS]; -uint32_t *h_found2[MAX_GPUS]; extern void quark_blake512_cpu_init(int thr_id, uint32_t threads); extern void quark_blake512_cpu_setBlock_80(void *pdata); @@ -143,8 +142,8 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, { const uint32_t first_nonce = pdata[19]; - int intensity = 256 * 256 * 20; - if (device_sm[device_map[thr_id]] == 520) intensity = 256 * 256 * 32; + int intensity = 256 * 256 * 10; + if (device_sm[device_map[thr_id]] == 520) intensity = 256 * 256 * 21*(1); uint32_t throughput = device_intensity(thr_id, __func__, intensity); // 19=256*256*8; throughput = min(throughput, (max_nonce - first_nonce)); @@ -166,7 +165,6 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, } CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], 64 * throughput), 0); // why 64 ? CUDA_CALL_OR_RET_X(cudaMallocHost(&(h_found[thr_id]), 4 * sizeof(uint32_t)), 0); - CUDA_CALL_OR_RET_X(cudaMallocHost(&(h_found2[thr_id]), 4 * sizeof(uint32_t)), 0); cuda_check_cpu_init(thr_id, throughput); init[thr_id] = true; } @@ -178,23 +176,19 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, do { int order = 0; - quark_blake512_cpu_hash_80(thr_id, throughput >> 1, pdata[19], d_hash[thr_id], order++); - quark_blake512_cpu_hash_80(thr_id, throughput >> 1, pdata[19] + (throughput / 2), d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - quark_bmw512_cpu_hash_64(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], order++); - quark_bmw512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - quark_groestl512_cpu_hash_64(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], order++); - quark_groestl512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - quark_skein512_cpu_hash_64(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], order++); - quark_skein512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - cuda_jh512Keccak512_cpu_hash_64(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], order++); - cuda_jh512Keccak512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - x11_luffaCubehash512_cpu_hash_64(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], order++); - x11_luffaCubehash512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - x11_shavite512_cpu_hash_64(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], order++); - x11_simd512_cpu_hash_64(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], order); - x11_shavite512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - x11_echo512_cpu_hash_64_final(thr_id, throughput >> 1, pdata[19], NULL, d_hash[thr_id], ptarget[7], h_found[thr_id], order++); + quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + quark_bmw512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + cuda_jh512Keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x11_luffaCubehash512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order); + cudaDeviceSynchronize(); + + x11_echo512_cpu_hash_64_final(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], ptarget[7], h_found[thr_id], order++); + h_found[thr_id][0] = 0xffffffff; if (h_found[thr_id][0] != 0xffffffff) { const uint32_t Htarg = ptarget[7]; @@ -226,40 +220,6 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, } } } - x11_simd512_cpu_hash_64(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), order++); - x11_echo512_cpu_hash_64_final(thr_id, throughput >> 1, pdata[19] + (throughput / 2), NULL, d_hash[thr_id] + (throughput * 32 / sizeof(int)), ptarget[7], h_found2[thr_id], order++); - if (h_found2[thr_id][0] != 0xffffffff) - { - const uint32_t Htarg = ptarget[7]; - uint32_t vhash64[8]; - be32enc(&endiandata[19], h_found2[thr_id][0]); - x11hash(vhash64, endiandata); - - if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) - { - int res = 1; - *hashes_done = pdata[19] - first_nonce + throughput; - if (h_found2[thr_id][1] != 0xffffffff) - { - pdata[21] = h_found2[thr_id][1]; - res++; - if (opt_benchmark) - applog(LOG_INFO, "GPU #%d Found second nounce %08x", thr_id, h_found2[thr_id][1], vhash64[7], Htarg); - } - pdata[19] = h_found2[thr_id][0]; - if (opt_benchmark) - applog(LOG_INFO, "GPU #%d Found nounce %08x", thr_id, h_found2[thr_id][0], vhash64[7], Htarg); - return res; - } - else - { - if (vhash64[7] != Htarg) - { - applog(LOG_INFO, "GPU #%d: result for %08x does not validate on CPU!", thr_id, h_found2[thr_id][0]); - } - } - } - pdata[19] += throughput; } while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > ((uint64_t)(pdata[19]) + (uint64_t)throughput))); From f6293a767eb331af938fa19f42f196608e860dcf Mon Sep 17 00:00:00 2001 From: KlausT Date: Thu, 5 Feb 2015 18:32:32 +0100 Subject: [PATCH 13/23] fix Animecoin bug --- quark/animecoin.cu | 22 ++++++++++++++-------- quark/cuda_bmw512.cu | 10 +++++----- 2 files changed, 19 insertions(+), 13 deletions(-) diff --git a/quark/animecoin.cu b/quark/animecoin.cu index 96ab889068..088641d7bb 100644 --- a/quark/animecoin.cu +++ b/quark/animecoin.cu @@ -167,7 +167,7 @@ extern "C" int scanhash_anime(int thr_id, uint32_t *pdata, unsigned long *hashes_done) { const uint32_t first_nonce = pdata[19]; - uint32_t throughput = device_intensity(thr_id, __func__, 1 << 19); // 256*256*8 + uint32_t throughput = device_intensity(thr_id, __func__, 1 << 20); // 256*256*8 throughput = min(throughput, (max_nonce - first_nonce)); if (opt_benchmark) @@ -258,31 +258,37 @@ extern "C" int scanhash_anime(int thr_id, uint32_t *pdata, if (foundnonces[0] != 0xffffffff) { const uint32_t Htarg = ptarget[7]; -/* uint32_t vhash64[8]; + uint32_t vhash64[8]; be32enc(&endiandata[19], foundnonces[0]); animehash(vhash64, endiandata); if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) -*/ { + { int res = 1; *hashes_done = pdata[19] - first_nonce + throughput; // check if there was some other ones... if (foundnonces[1] != 0xffffffff) { - pdata[21] = foundnonces[1]; - res++; - if (opt_benchmark) applog(LOG_INFO, "GPU #%d: Found second nonce $%08X", thr_id, foundnonces[1]); + be32enc(&endiandata[19], foundnonces[1]); + animehash(vhash64, endiandata); + if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) + { + + pdata[21] = foundnonces[1]; + res++; + if (opt_benchmark) applog(LOG_INFO, "GPU #%d: Found second nonce $%08X", thr_id, foundnonces[1]); + } } pdata[19] = foundnonces[0]; if (opt_benchmark) applog(LOG_INFO, "GPU #%d: Found nonce $%08X", thr_id, foundnonces[0]); return res; } -/* else + else { if (vhash64[7] != Htarg) // don't show message if it is equal but fails fulltest applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU!", thr_id, foundnonces[0]); } -*/ } + } pdata[19] += throughput; } while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > ((uint64_t)(pdata[19]) + (uint64_t)throughput))); diff --git a/quark/cuda_bmw512.cu b/quark/cuda_bmw512.cu index a425e9667d..68477ecc8f 100644 --- a/quark/cuda_bmw512.cu +++ b/quark/cuda_bmw512.cu @@ -335,8 +335,8 @@ void quark_bmw512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t * { 0xACADAEAFUL, 0xA8A9AAABUL }, { 0xB4B5B6B7UL, 0xB0B1B2B3UL }, { 0xBCBDBEBFUL, 0xB8B9BABBUL }, - { 0xC4C5C6C7UL, 0xC0C1C2C3UL, }, - { 0xCCCDCECFUL, 0xC8C9CACBUL, }, + { 0xC4C5C6C7UL, 0xC0C1C2C3UL }, + { 0xCCCDCECFUL, 0xC8C9CACBUL }, { 0xD4D5D6D7UL, 0xD0D1D2D3UL }, { 0xDCDDDEDFUL, 0xD8D9DADBUL }, { 0xE4E5E6E7UL, 0xE0E1E2E3UL }, @@ -398,8 +398,8 @@ void quark_bmw512_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint64_t * { 0xACADAEAFUL, 0xA8A9AAABUL }, { 0xB4B5B6B7UL, 0xB0B1B2B3UL }, { 0xBCBDBEBFUL, 0xB8B9BABBUL }, - { 0xC4C5C6C7UL, 0xC0C1C2C3UL, }, - { 0xCCCDCECFUL, 0xC8C9CACBUL, }, + { 0xC4C5C6C7UL, 0xC0C1C2C3UL }, + { 0xCCCDCECFUL, 0xC8C9CACBUL }, { 0xD4D5D6D7UL, 0xD0D1D2D3UL }, { 0xDCDDDEDFUL, 0xD8D9DADBUL }, { 0xE4E5E6E7UL, 0xE0E1E2E3UL }, @@ -415,7 +415,7 @@ void quark_bmw512_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint64_t * message[i] = vectorize(c_PaddedMessage80[i]); // die Nounce durch die thread-spezifische ersetzen - message[9].x = cuda_swab32(nounce); //REPLACE_HIWORD(message[9], cuda_swab32(nounce)); + message[9].y = cuda_swab32(nounce); //REPLACE_HIWORD(message[9], cuda_swab32(nounce)); // Compression 1 Compression512(message, h); From 87e1f5767829a7b7bb826932dc85f02f7c7345ac Mon Sep 17 00:00:00 2001 From: sp-hash Date: Thu, 5 Feb 2015 15:23:00 +0100 Subject: [PATCH 14/23] fixed x11 --- x11/x11.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/x11/x11.cu b/x11/x11.cu index fc0c0eeaf8..6cff75b8a8 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -188,7 +188,7 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, cudaDeviceSynchronize(); x11_echo512_cpu_hash_64_final(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], ptarget[7], h_found[thr_id], order++); - h_found[thr_id][0] = 0xffffffff; +// h_found[thr_id][0] = 0xffffffff; if (h_found[thr_id][0] != 0xffffffff) { const uint32_t Htarg = ptarget[7]; From 5829510434742399090e1533dbf0aa85ee51ad7c Mon Sep 17 00:00:00 2001 From: sp-hash Date: Thu, 5 Feb 2015 18:27:45 +0100 Subject: [PATCH 15/23] Bether launchconfig on the 750ti (echo) --- x11/cuda_x11_echo.cu | 31 ++++++++++++++++++++++--------- 1 file changed, 22 insertions(+), 9 deletions(-) diff --git a/x11/cuda_x11_echo.cu b/x11/cuda_x11_echo.cu index 7fee109512..ec98adc9aa 100644 --- a/x11/cuda_x11_echo.cu +++ b/x11/cuda_x11_echo.cu @@ -309,13 +309,12 @@ __device__ __forceinline__ void cuda_echo_round( hash[i] ^= W[i]; } - - +/* __device__ __forceinline__ -void echo_gpu_init(uint32_t *const __restrict__ sharedMemory) +void echo_gpu_init_128(uint32_t *const __restrict__ sharedMemory) { - /* each thread startup will fill a uint32 */ - if (threadIdx.x < 128) { + if (threadIdx.x < 128) + { sharedMemory[threadIdx.x] = d_AES0[threadIdx.x]; sharedMemory[threadIdx.x + 256] = d_AES1[threadIdx.x]; sharedMemory[threadIdx.x + 512] = d_AES2[threadIdx.x]; @@ -327,9 +326,23 @@ void echo_gpu_init(uint32_t *const __restrict__ sharedMemory) sharedMemory[threadIdx.x + 64 * 2 + 768] = d_AES3[threadIdx.x + 64 * 2]; } } +*/ + + +__device__ __forceinline__ +void echo_gpu_init(uint32_t *const __restrict__ sharedMemory) +{ + /* each thread startup will fill a uint32 */ + if (threadIdx.x < 128) { + sharedMemory[threadIdx.x] = d_AES0[threadIdx.x]; + sharedMemory[threadIdx.x + 256] = d_AES1[threadIdx.x]; + sharedMemory[threadIdx.x + 512] = d_AES2[threadIdx.x]; + sharedMemory[threadIdx.x + 768] = d_AES3[threadIdx.x]; + } +} -__global__ __launch_bounds__(128) +__global__ __launch_bounds__(256, 4) void x11_echo512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t *const __restrict__ g_hash, const uint32_t *const __restrict__ g_nonceVector) { __shared__ uint32_t sharedMemory[1024]; @@ -355,7 +368,7 @@ __host__ void x11_echo512_cpu_init(int thr_id, uint32_t threads) __host__ void x11_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) { - const uint32_t threadsperblock = 128; + const uint32_t threadsperblock = 256; // berechne wie viele Thread Blocks wir brauchen dim3 grid((threads + threadsperblock-1)/threadsperblock); @@ -370,7 +383,7 @@ __host__ void x11_echo512_cpu_free(int32_t thr_id) cudaFreeHost(&d_nonce[thr_id]); } -__global__ __launch_bounds__(128) +__global__ __launch_bounds__(256, 4) void x11_echo512_gpu_hash_64_final(uint32_t threads, uint32_t startNounce, uint64_t *const __restrict__ g_hash, const uint32_t *const __restrict__ g_nonceVector, uint32_t *const __restrict__ d_found, uint32_t target) { __shared__ uint32_t sharedMemory[1024]; @@ -698,7 +711,7 @@ void x11_echo512_gpu_hash_64_final(uint32_t threads, uint32_t startNounce, uint6 } __host__ void x11_echo512_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, uint32_t target, uint32_t *h_found, int order) { - const uint32_t threadsperblock = 128; + const uint32_t threadsperblock = 256; // berechne wie viele Thread Blocks wir brauchen dim3 grid((threads + threadsperblock - 1) / threadsperblock); From e0f7eed32cb9d4e27b64b7c16a6b2df97d99a25c Mon Sep 17 00:00:00 2001 From: KlausT Date: Thu, 5 Feb 2015 20:07:26 +0100 Subject: [PATCH 16/23] remove cudaDeviceSynchronize --- x11/x11.cu | 3 --- 1 file changed, 3 deletions(-) diff --git a/x11/x11.cu b/x11/x11.cu index 6cff75b8a8..04b88572e2 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -185,10 +185,7 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order); - cudaDeviceSynchronize(); - x11_echo512_cpu_hash_64_final(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], ptarget[7], h_found[thr_id], order++); -// h_found[thr_id][0] = 0xffffffff; if (h_found[thr_id][0] != 0xffffffff) { const uint32_t Htarg = ptarget[7]; From 77786ec91ab135c475c86652e58dc3643762c2ba Mon Sep 17 00:00:00 2001 From: KlausT Date: Thu, 5 Feb 2015 21:01:28 +0100 Subject: [PATCH 17/23] add error checks --- Algo256/blake256.cu | 10 ++++------ Algo256/cuda_fugue256.cu | 2 +- Algo256/cuda_keccak256.cu | 2 +- cuda_checkhash.cu | 4 ++-- cuda_groestlcoin.cu | 2 +- cuda_myriadgroestl.cu | 2 +- pentablake.cu | 2 +- quark/cuda_skein512.cu | 2 +- x15/cuda_x15_whirlpool.cu | 2 +- 9 files changed, 13 insertions(+), 15 deletions(-) diff --git a/Algo256/blake256.cu b/Algo256/blake256.cu index 8a3a9f1e23..09f5dff106 100644 --- a/Algo256/blake256.cu +++ b/Algo256/blake256.cu @@ -339,12 +339,10 @@ static uint32_t blake256_cpu_hash_16(const int thr_id, const uint32_t threads, c return result; blake256_gpu_hash_16 <<>> (threads, startNonce, d_resNonce[thr_id], highTarget, (int) rounds, opt_tracegpu); - if (cudaSuccess == cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) { - //cudaDeviceSynchronize(); /* 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]; - } + CUDA_SAFE_CALL(cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)); + 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; } diff --git a/Algo256/cuda_fugue256.cu b/Algo256/cuda_fugue256.cu index 08950a5d41..e18bf412ef 100644 --- a/Algo256/cuda_fugue256.cu +++ b/Algo256/cuda_fugue256.cu @@ -765,5 +765,5 @@ __host__ void fugue256_cpu_hash(int thr_id, uint32_t threads, int startNounce, v fugue256_gpu_hash<<>>(thr_id, threads, startNounce, d_fugue256_hashoutput[thr_id], d_resultNonce[thr_id]); //cudaMemcpy(outputHashes, d_fugue256_hashoutput[thr_id], 8 * sizeof(uint32_t), cudaMemcpyDeviceToHost); - cudaMemcpy(nounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost); + CUDA_SAFE_CALL(cudaMemcpy(nounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost)); } diff --git a/Algo256/cuda_keccak256.cu b/Algo256/cuda_keccak256.cu index 7b99cd7fcc..ef63aff5ca 100644 --- a/Algo256/cuda_keccak256.cu +++ b/Algo256/cuda_keccak256.cu @@ -629,7 +629,7 @@ void keccak256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, u keccak256_gpu_hash_80<<>>(threads, startNounce, d_outputHash, d_KNonce[thr_id]); //MyStreamSynchronize(NULL, order, thr_id); - cudaMemcpy(h_nounce, d_KNonce[thr_id], 4 * sizeof(uint32_t), cudaMemcpyDeviceToHost); + CUDA_SAFE_CALL(cudaMemcpy(h_nounce, d_KNonce[thr_id], 4 * sizeof(uint32_t), cudaMemcpyDeviceToHost)); } __global__ __launch_bounds__(256,3) diff --git a/cuda_checkhash.cu b/cuda_checkhash.cu index 070c307b43..c1e6224570 100644 --- a/cuda_checkhash.cu +++ b/cuda_checkhash.cu @@ -90,7 +90,7 @@ void cuda_checkhash_64(uint32_t threads, uint32_t startNounce, uint32_t *hash, u __host__ uint32_t cuda_check_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash) { - cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t)); + CUDA_SAFE_CALL(cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t))); const uint32_t threadsperblock = 512; @@ -212,7 +212,7 @@ uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNoun __host__ void cuda_check_quarkcoin(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order, uint32_t *resNonces) { - cudaMemset(d_resNonces[thr_id], 0xff, 2*sizeof(uint32_t)); + CUDA_SAFE_CALL(cudaMemset(d_resNonces[thr_id], 0xff, 2*sizeof(uint32_t))); const uint32_t threadsperblock = 256; diff --git a/cuda_groestlcoin.cu b/cuda_groestlcoin.cu index 6baeb89b77..a2b79f2ab7 100644 --- a/cuda_groestlcoin.cu +++ b/cuda_groestlcoin.cu @@ -121,5 +121,5 @@ __host__ void groestlcoin_cpu_hash(int thr_id, uint32_t threads, uint32_t startN cudaMemset(d_resultNonce[thr_id], 0xFF, 2 * sizeof(uint32_t)); groestlcoin_gpu_hash_quad<<>>(threads, startNounce, d_resultNonce[thr_id]); - cudaMemcpy(nounce, d_resultNonce[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost); + CUDA_SAFE_CALL(cudaMemcpy(nounce, d_resultNonce[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost)); } diff --git a/cuda_myriadgroestl.cu b/cuda_myriadgroestl.cu index 4ba7f9555c..12233dd21e 100644 --- a/cuda_myriadgroestl.cu +++ b/cuda_myriadgroestl.cu @@ -326,5 +326,5 @@ __host__ void myriadgroestl_cpu_hash(int thr_id, uint32_t threads, uint32_t star myriadgroestl_gpu_hash_quad2<<>>(threads, startNounce, d_resultNonce[thr_id], d_outputHashes[thr_id]); - cudaMemcpy(nounce, d_resultNonce[thr_id], 4*sizeof(uint32_t), cudaMemcpyDeviceToHost); + CUDA_SAFE_CALL(cudaMemcpy(nounce, d_resultNonce[thr_id], 4*sizeof(uint32_t), cudaMemcpyDeviceToHost)); } diff --git a/pentablake.cu b/pentablake.cu index 3b184e53af..5aaee9c271 100644 --- a/pentablake.cu +++ b/pentablake.cu @@ -493,7 +493,7 @@ extern "C" int scanhash_pentablake(int thr_id, uint32_t *pdata, const uint32_t * pentablake_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], order++); pentablake_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], order++); pentablake_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], order++); - + CUDA_SAFE_CALL(cudaGetLastError()); uint32_t foundNonce = pentablake_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id], order++); if (foundNonce != UINT32_MAX) { diff --git a/quark/cuda_skein512.cu b/quark/cuda_skein512.cu index ac0be31667..1f2ba25f2d 100644 --- a/quark/cuda_skein512.cu +++ b/quark/cuda_skein512.cu @@ -536,6 +536,6 @@ void quark_skein512_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t sta cudaMemset(d_nonce[thr_id], 0xff, 2*sizeof(uint32_t)); quark_skein512_gpu_hash_64_final<< > >(threads, startNounce, (uint64_t*)d_hash, d_nonceVector, d_nonce[thr_id], target); - cudaMemcpy(h_nonce, d_nonce[thr_id], 2*sizeof(uint32_t), cudaMemcpyDeviceToHost); + CUDA_SAFE_CALL(cudaMemcpy(h_nonce, d_nonce[thr_id], 2*sizeof(uint32_t), cudaMemcpyDeviceToHost)); } diff --git a/x15/cuda_x15_whirlpool.cu b/x15/cuda_x15_whirlpool.cu index 4729d9b902..1371e09c3f 100644 --- a/x15/cuda_x15_whirlpool.cu +++ b/x15/cuda_x15_whirlpool.cu @@ -2609,7 +2609,7 @@ extern uint32_t whirlpool512_cpu_finalhash_64(int thr_id, uint32_t threads, uint oldwhirlpool_gpu_finalhash_64<<>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector,d_WNonce[thr_id]); //MyStreamSynchronize(NULL, order, thr_id); - cudaMemcpy(d_wnounce[thr_id], d_WNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost); + CUDA_SAFE_CALL(cudaMemcpy(d_wnounce[thr_id], d_WNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost)); result = *d_wnounce[thr_id]; return result; } From 28f8a62e43f94139b15dff6cf6e670745a5a492d Mon Sep 17 00:00:00 2001 From: KlausT Date: Thu, 5 Feb 2015 21:21:01 +0100 Subject: [PATCH 18/23] add device_props[] variable, to be used in a later commit --- cuda.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cuda.cpp b/cuda.cpp index ec31727fd4..e3ddd0576a 100644 --- a/cuda.cpp +++ b/cuda.cpp @@ -20,6 +20,8 @@ #include "cuda_runtime.h" +cudaDeviceProp device_props[MAX_GPUS]; + // CUDA Devices on the System int cuda_num_devices() { @@ -64,6 +66,7 @@ void cuda_devicenames() cudaDeviceProp props; cudaGetDeviceProperties(&props, device_map[i]); + device_props[i] = props; device_name[i] = strdup(props.name); device_sm[i] = (props.major * 100 + props.minor * 10); } From 170ec05d4e656f11babc4372ff28b722ce4fba2e Mon Sep 17 00:00:00 2001 From: KlausT Date: Thu, 5 Feb 2015 23:07:16 +0100 Subject: [PATCH 19/23] delete stuff that is defined twice --- x11/cuda_x11_luffa512_Cubehash.cu | 56 +++++++++++++------------------ 1 file changed, 23 insertions(+), 33 deletions(-) diff --git a/x11/cuda_x11_luffa512_Cubehash.cu b/x11/cuda_x11_luffa512_Cubehash.cu index ef5bb7f963..45b7dd50f8 100644 --- a/x11/cuda_x11_luffa512_Cubehash.cu +++ b/x11/cuda_x11_luffa512_Cubehash.cu @@ -44,6 +44,14 @@ typedef struct { #define LROT(x, bits) __funnelshift_l(x, x, bits) #endif +#define CUBEHASH_ROUNDS 16 /* this is r for CubeHashr/b */ +#define CUBEHASH_BLOCKBYTES 32 /* this is b for CubeHashr/b */ + +#define ROTATEUPWARDS7(a) LROT(a,7) +#define ROTATEUPWARDS11(a) LROT(a,11) + +#define SWAP(a,b) { uint32_t u = a; a = b; b = u; } + #define TWEAK(a0,a1,a2,a3,j)\ a0 = LROT(a0,j);\ a1 = LROT(a1,j);\ @@ -121,6 +129,21 @@ __device__ __constant__ uint32_t c_IV[40] = { 0xf5df3999,0x0fc688f1,0xb07224cc,0x03e86cea}; */ +__device__ __constant__ +static const uint32_t c_IV_512[32] = { + + 0x2AEA2A61, 0x50F494D4, 0x2D538B8B, + 0x4167D83E, 0x3FEE2313, 0xC701CF8C, + 0xCC39968E, 0x50AC5695, 0x4D42C787, + 0xA647A8B3, 0x97CF0BEF, 0x825B4537, + 0xEEF864D2, 0xF22090C4, 0xD0E5CD33, + 0xA23911AE, 0xFCD398D9, 0x148FE485, + 0x1B017BEF, 0xB6444532, 0x6A536159, + 0x2FF5781C, 0x91FA7934, 0x0DBADEA9, + 0xD65C8A2B, 0xA5A70E75, 0xB1C62456, + 0xBC796576, 0x1921C8F7, 0xE7989AF1, + 0x7795D246, 0xD43E3B44 +}; __device__ __constant__ uint32_t c_CNS[80] = { 0x303994a6,0xe0337818,0xc0e65299,0x441ba90d, @@ -566,39 +589,6 @@ void finalization512(hashState *state, uint32_t *b) } } - -typedef unsigned char BitSequence; - -#define CUBEHASH_ROUNDS 16 /* this is r for CubeHashr/b */ -#define CUBEHASH_BLOCKBYTES 32 /* this is b for CubeHashr/b */ - -#if __CUDA_ARCH__ < 350 -#define LROT(x,bits) ((x << bits) | (x >> (32 - bits))) -#else -#define LROT(x, bits) __funnelshift_l(x, x, bits) -#endif - -#define ROTATEUPWARDS7(a) LROT(a,7) -#define ROTATEUPWARDS11(a) LROT(a,11) - -#define SWAP(a,b) { uint32_t u = a; a = b; b = u; } - -__device__ __constant__ -static const uint32_t c_IV_512[32] = { - - 0x2AEA2A61, 0x50F494D4, 0x2D538B8B, - 0x4167D83E, 0x3FEE2313, 0xC701CF8C, - 0xCC39968E, 0x50AC5695, 0x4D42C787, - 0xA647A8B3, 0x97CF0BEF, 0x825B4537, - 0xEEF864D2, 0xF22090C4, 0xD0E5CD33, - 0xA23911AE, 0xFCD398D9, 0x148FE485, - 0x1B017BEF, 0xB6444532, 0x6A536159, - 0x2FF5781C, 0x91FA7934, 0x0DBADEA9, - 0xD65C8A2B, 0xA5A70E75, 0xB1C62456, - 0xBC796576, 0x1921C8F7, 0xE7989AF1, - 0x7795D246, 0xD43E3B44 -}; - __device__ __forceinline__ void rrounds(uint32_t x[2][2][2][2][2]) { int r; From ebce5183da5b530865f812b2a4df50463d4502d1 Mon Sep 17 00:00:00 2001 From: KlausT Date: Thu, 5 Feb 2015 23:19:52 +0100 Subject: [PATCH 20/23] X11: verify second nonce --- x11/x11.cu | 22 ++++++++++++++++++---- 1 file changed, 18 insertions(+), 4 deletions(-) diff --git a/x11/x11.cu b/x11/x11.cu index 04b88572e2..cda20e5f7d 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -199,10 +199,24 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, *hashes_done = pdata[19] - first_nonce + throughput; if (h_found[thr_id][1] != 0xffffffff) { - pdata[21] = h_found[thr_id][1]; - res++; - if (opt_benchmark) - applog(LOG_INFO, "GPU #%d Found second nounce %08x", thr_id, h_found[thr_id][1], vhash64[7], Htarg); + be32enc(&endiandata[19], h_found[thr_id][1]); + x11hash(vhash64, endiandata); + if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) + { + + pdata[21] = h_found[thr_id][1]; + res++; + if (opt_benchmark) + applog(LOG_INFO, "GPU #%d Found second nounce %08x", thr_id, h_found[thr_id][1], vhash64[7], Htarg); + } + else + { + if (vhash64[7] != Htarg) + { + applog(LOG_INFO, "GPU #%d: result for %08x does not validate on CPU!", thr_id, h_found[thr_id][1]); + } + } + } pdata[19] = h_found[thr_id][0]; if (opt_benchmark) From 0e09573a38f996bd382555a8d3561ed4fda62556 Mon Sep 17 00:00:00 2001 From: KlausT Date: Fri, 6 Feb 2015 00:21:24 +0100 Subject: [PATCH 21/23] Quark: verify second nonce --- quark/quarkcoin.cu | 17 ++++++++++++++--- 1 file changed, 14 insertions(+), 3 deletions(-) diff --git a/quark/quarkcoin.cu b/quark/quarkcoin.cu index bbff75bd99..3532a151bb 100644 --- a/quark/quarkcoin.cu +++ b/quark/quarkcoin.cu @@ -251,9 +251,20 @@ extern "C" int scanhash_quark(int thr_id, uint32_t *pdata, // check if there was some other ones... if (foundnonces[1] != 0xffffffff) { - pdata[21] = foundnonces[1]; - res++; - if (opt_benchmark) applog(LOG_INFO, "GPU #%d: Found second nonce $%08X", thr_id, foundnonces[1]); + be32enc(&endiandata[19], foundnonces[1]); + quarkhash(vhash64, endiandata); + + if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) + { + pdata[21] = foundnonces[1]; + res++; + if (opt_benchmark) applog(LOG_INFO, "GPU #%d: Found second nonce $%08X", thr_id, foundnonces[1]); + } + else + { + if (vhash64[7] != Htarg) // don't show message if it is equal but fails fulltest + applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU!", thr_id, foundnonces[1]); + } } pdata[19] = foundnonces[0]; if (opt_benchmark) applog(LOG_INFO, "GPU #%d: Found nonce $%08X", thr_id, foundnonces[0]); From aeccf8a32855ecd8fedad10e9fcd6d7f62e8d814 Mon Sep 17 00:00:00 2001 From: KlausT Date: Fri, 6 Feb 2015 01:01:29 +0100 Subject: [PATCH 22/23] small shavite optimization --- x11/cuda_x11_shavite512.cu | 16 +++++----------- 1 file changed, 5 insertions(+), 11 deletions(-) diff --git a/x11/cuda_x11_shavite512.cu b/x11/cuda_x11_shavite512.cu index d531da9321..d67b58e9fc 100644 --- a/x11/cuda_x11_shavite512.cu +++ b/x11/cuda_x11_shavite512.cu @@ -1350,20 +1350,14 @@ void x11_shavite512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t pE = state[0xE]; pF = state[0xF]; - x0 = p4; - x1 = p5; - x2 = p6; - x3 = p7; - - rk[0] = msg[0]; - x0 ^= msg[0]; + x0 = msg[0] ^ state[0x4]; rk[1] = msg[1]; - x1 ^= msg[1]; + x1 = msg[1] ^ state[0x5]; rk[2] = msg[2]; - x2 ^= msg[2]; + x2 = msg[2] ^ state[0x6]; rk[3] = msg[3]; - x3 ^= msg[3]; + x3 = msg[3] ^ state[0x7]; AES_ROUND_NOKEY(sharedMemory, x0, x1, x2, x3); rk[4] = msg[4]; @@ -2592,7 +2586,7 @@ void x11_shavite512_gpu_hash_80(uint32_t threads, uint32_t startNounce, void *ou uint32_t msg[32]; #pragma unroll 32 - for(int i=0;i<32;i++) { + for(int i=0;i<31;i++) { msg[i] = c_PaddedMessage80[i]; } msg[19] = cuda_swab32(nounce); From 7306a68fea9e7f0ad8f3e6fb4afdf038c256474b Mon Sep 17 00:00:00 2001 From: KlausT Date: Fri, 6 Feb 2015 04:19:29 +0100 Subject: [PATCH 23/23] echo bug fix --- x11/cuda_x11_echo.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/x11/cuda_x11_echo.cu b/x11/cuda_x11_echo.cu index ec98adc9aa..6591901201 100644 --- a/x11/cuda_x11_echo.cu +++ b/x11/cuda_x11_echo.cu @@ -333,7 +333,7 @@ __device__ __forceinline__ void echo_gpu_init(uint32_t *const __restrict__ sharedMemory) { /* each thread startup will fill a uint32 */ - if (threadIdx.x < 128) { + if (threadIdx.x < 256) { sharedMemory[threadIdx.x] = d_AES0[threadIdx.x]; sharedMemory[threadIdx.x + 256] = d_AES1[threadIdx.x]; sharedMemory[threadIdx.x + 512] = d_AES2[threadIdx.x];