diff options
author | fireice-uk <fireice-uk@users.noreply.github.com> | 2017-10-17 21:30:29 +0100 |
---|---|---|
committer | fireice-uk <fireice-uk@users.noreply.github.com> | 2017-10-22 13:15:18 +0100 |
commit | d3c5911f8e486ac0d08c5160f6d9bd934c697b91 (patch) | |
tree | 66441e722f2bbc8b99e6ba597b07bd6bb38aec59 /xmrstak | |
parent | 27b839b8d2275c0e78a07cb03d27afb96ed6c7b9 (diff) | |
download | xmr-stak-d3c5911f8e486ac0d08c5160f6d9bd934c697b91.zip xmr-stak-d3c5911f8e486ac0d08c5160f6d9bd934c697b91.tar.gz |
move nicehash check into nonce calc
Diffstat (limited to 'xmrstak')
-rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 6 | ||||
-rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.hpp | 2 | ||||
-rw-r--r-- | xmrstak/backend/amd/minethd.cpp | 5 | ||||
-rw-r--r-- | xmrstak/backend/cpu/minethd.cpp | 16 | ||||
-rw-r--r-- | xmrstak/backend/globalStates.hpp | 11 | ||||
-rw-r--r-- | xmrstak/backend/nvidia/minethd.cpp | 5 |
6 files changed, 22 insertions, 23 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index f9908cb..37adc5b 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -844,7 +844,8 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput) } }*/ - if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1], 1, &ctx->Nonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS) + size_t tmpNonce = ctx->Nonce; + if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS) { printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 1); return ERR_OCL_API; @@ -897,7 +898,8 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput) BranchNonces[i] = ((BranchNonces[i] + w_size - 1u) / w_size) * w_size; // number of global threads must be a multiple of the work group size (w_size) assert(BranchNonces[i]%w_size == 0); - if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[i + 3], 1, &ctx->Nonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS) + size_t tmpNonce = ctx->Nonce; + if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[i + 3], 1, &tmpNonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS) { printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), i + 3); return ERR_OCL_API; diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index 5ff7ea1..c2d708d 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -36,7 +36,7 @@ struct GpuContext int computeUnits; std::string name; - size_t Nonce; + uint32_t Nonce; }; diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 8658e5c..ca0b110 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -200,10 +200,7 @@ void minethd::work_main() //Allocate a new nonce every 16 rounds if((round_ctr++ & 0xF) == 0) { - if(oWork.bNiceHash) - pGpuCtx->Nonce = globalStates::inst().calc_start_nonce(pGpuCtx->Nonce & 0xFF000000u, h_per_round * 16); - else - pGpuCtx->Nonce = globalStates::inst().calc_start_nonce(0, h_per_round * 16); + globalStates::inst().calc_start_nonce(pGpuCtx->Nonce, oWork.bNiceHash, h_per_round * 16); } cl_uint results[0x100]; diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 542e999..ac80cc2 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -348,6 +348,9 @@ void minethd::work_main() assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID)); memcpy(result.sJobID, oWork.sJobID, sizeof(job_result::sJobID)); + if(oWork.bNiceHash) + result.iNonce = *piNonce; + while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { if ((iCount++ & 0xF) == 0) //Store stats every 16 hashes @@ -360,10 +363,7 @@ void minethd::work_main() if((nonce_ctr++ & (nonce_chunk-1)) == 0) { - if(oWork.bNiceHash) - result.iNonce = globalStates::inst().calc_start_nonce(*piNonce & 0xFF000000, nonce_chunk); - else - result.iNonce = globalStates::inst().calc_start_nonce(0, nonce_chunk); + globalStates::inst().calc_start_nonce(result.iNonce, oWork.bNiceHash, nonce_chunk); } *piNonce = ++result.iNonce; @@ -466,6 +466,9 @@ void minethd::double_work_main() assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID)); + if(oWork.bNiceHash) + iNonce = *piNonce0; + while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { if ((iCount & 0x7) == 0) //Store stats every 16 hashes @@ -480,10 +483,7 @@ void minethd::double_work_main() if((nonce_ctr++ & (nonce_chunk/2 - 1)) == 0) { - if(oWork.bNiceHash) - iNonce = globalStates::inst().calc_start_nonce(*piNonce0 & 0xFF000000, nonce_chunk); - else - iNonce = globalStates::inst().calc_start_nonce(0, nonce_chunk); + globalStates::inst().calc_start_nonce(iNonce, oWork.bNiceHash, nonce_chunk); } diff --git a/xmrstak/backend/globalStates.hpp b/xmrstak/backend/globalStates.hpp index 855488d..50dc753 100644 --- a/xmrstak/backend/globalStates.hpp +++ b/xmrstak/backend/globalStates.hpp @@ -33,11 +33,14 @@ struct globalStates //pool_data is in-out winapi style void switch_work(miner_work& pWork, pool_data& dat); - inline uint32_t calc_start_nonce(uint32_t nicehash_nonce, uint32_t reserve_count) + inline void calc_start_nonce(uint32_t& nonce, bool use_nicehash, uint32_t reserve_count) { - uint32_t debug_nonce = nicehash_nonce | iGlobalNonce.fetch_add(reserve_count); - printer::inst()->print_msg(L1, "DEBUG: start_nonce assigned nh: %.8x rc: %.8x nonce: %.8x", nicehash_nonce, reserve_count, debug_nonce); - return debug_nonce; + if(use_nicehash) + nonce = (nonce & 0xFF000000) | iGlobalNonce.fetch_add(reserve_count); + else + nonce = iGlobalNonce.fetch_add(reserve_count); + + printer::inst()->print_msg(L1, "DEBUG: start_nonce assigned rc: %.8x nonce: %.8x", reserve_count, nonce); } miner_work oGlobalWork; diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index f82d56b..20d578d 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -232,10 +232,7 @@ void minethd::work_main() //Allocate a new nonce every 16 rounds if((round_ctr++ & 0xF) == 0) { - if(oWork.bNiceHash) - iNonce = globalStates::inst().calc_start_nonce(iNonce & 0xFF000000u, h_per_round * 16); - else - iNonce = globalStates::inst().calc_start_nonce(0, h_per_round * 16); + iNonce = globalStates::inst().calc_start_nonce(iNonce, oWork.bNiceHash, h_per_round * 16); } uint32_t foundNonce[10]; |