move nicehash check into nonce calc

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;

xmrstak/backend/amd/amd_gpu/gpu.hpp

@@ -36,7 +36,7 @@ struct GpuContext
 	int computeUnits;
 	std::string name;
 
-	size_t Nonce;
+	uint32_t Nonce;
 
 };
 

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];

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);
 			}
 
 

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;

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];