From d3c5911f8e486ac0d08c5160f6d9bd934c697b91 Mon Sep 17 00:00:00 2001
From: fireice-uk <fireice-uk@users.noreply.github.com>
Date: Tue, 17 Oct 2017 21:30:29 +0100
Subject: [PATCH] move nicehash check into nonce calc
---
xmrstak/backend/amd/amd_gpu/gpu.cpp | 6 ++++--
xmrstak/backend/amd/amd_gpu/gpu.hpp | 2 +-
xmrstak/backend/amd/minethd.cpp | 5 +----
xmrstak/backend/cpu/minethd.cpp | 16 ++++++++--------
xmrstak/backend/globalStates.hpp | 11 +++++++----
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];
--
GitLab