diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index f9908cb0e6b97f9f2287dc85f8d1b4c6501d938f..37adc5b3775f24ae19d6970374a2c16739473b94 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 5ff7ea17bb01d1288db2b3c49076aedda9872c58..c2d708d9032b726b37fcb016cedac3045140a32a 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 8658e5cf3120cbf7c4f18970b60012bc5f2ef778..ca0b110b7475eae1d315af578c668eab4054027d 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 542e99978635f8ca7e21efc1a22a1dd2787c284a..ac80cc20a49301372a23a306a8be562532345165 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 855488ddedc8c20f286ef99c011a3e25d76b0381..50dc753626de42ca23808355cada01f820af8dc6 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 f82d56bb4da3c1d0b4bbea9854a0762bc47b4bb0..20d578de547265a7aeeb35565a93ed8cb4394816 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];