diff --git a/Algo256/blake256.cu b/Algo256/blake256.cu index 21b0276eae..09f5dff106 100644 --- a/Algo256/blake256.cu +++ b/Algo256/blake256.cu @@ -339,13 +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); - 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]; - 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/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/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/bitcoin.cu b/bitcoin.cu index 4eabd83b3f..97b537bbcf 100644 --- a/bitcoin.cu +++ b/bitcoin.cu @@ -112,13 +112,13 @@ void bitcoin_hash(uint32_t *output, const uint32_t *data, uint32_t nonce, const b = a; a = t1 + t2; } - 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[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]); } @@ -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/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/ccminer.cpp b/ccminer.cpp index f93d37793c..9fe552cbbb 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -222,7 +222,7 @@ Usage: " PROGRAM_NAME " [OPTIONS]\n\ Options:\n\ -a, --algo=ALGO specify the hash algorithm to use\n\ anime Animecoin\n\ - bitcoin use to mine Bitcoin\n\ + bitcoin Bitcoin\n\ blake Blake 256 (SFR/NEOS)\n\ blakecoin Fast Blake 256 (8 rounds)\n\ deep Deepcoin\n\ diff --git a/cuda.cpp b/cuda.cpp index 95c8221bdf..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() { @@ -32,7 +34,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); @@ -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); } 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 ce8ea18a5b..a2b79f2ab7 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; + } + } } } @@ -92,9 +71,12 @@ 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], 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 +118,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); + CUDA_SAFE_CALL(cudaMemcpy(nounce, d_resultNonce[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost)); } diff --git a/cuda_helper.h b/cuda_helper.h index b3834c24ec..955b6c75f3 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; } @@ -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)); } diff --git a/cuda_myriadgroestl.cu b/cuda_myriadgroestl.cu index d404afe7b4..12233dd21e 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); @@ -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/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/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; diff --git a/lyra2/lyra2RE.cu b/lyra2/lyra2RE.cu index 0453b3a96b..b3c2ae9873 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/pentablake.cu b/pentablake.cu index 1a2a9cde50..ba3fc87b8d 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/quark/quarkcoin.cu b/quark/quarkcoin.cu index aeb64af244..8a05f25506 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); @@ -247,9 +247,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]); 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/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]; 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; 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); diff --git a/x11/fresh.cu b/x11/fresh.cu index 8fb50a2d75..b28a07b429 100644 --- a/x11/fresh.cu +++ b/x11/fresh.cu @@ -86,7 +86,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 1112573dea..61653b823a 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -145,6 +145,7 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, 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)); if (opt_benchmark) @@ -152,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); @@ -197,10 +198,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) diff --git a/x13/x13.cu b/x13/x13.cu index 10f953f3c8..7e99742af3 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/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; } 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);