From f13f70c28eb9e4b38d3b4932f5845f1d5cc01906 Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Fri, 13 Oct 2017 16:35:10 +0100 Subject: Implement pool-controlled nonce allocation --- xmrstak/backend/amd/minethd.cpp | 19 ++++++++++++------- 1 file changed, 12 insertions(+), 7 deletions(-) (limited to 'xmrstak/backend/amd') diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 5ca10d3..048c3f0 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -167,13 +167,11 @@ void minethd::consume_work() void minethd::work_main() { uint64_t iCount = 0; - cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/); - + pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39); globalStates::inst().iConsumeCnt++; - uint32_t* piNonce = (uint32_t*)(oWork.bWorkBlob + 39); while (bQuit == 0) { @@ -190,10 +188,8 @@ void minethd::work_main() continue; } - if(oWork.bNiceHash) - pGpuCtx->Nonce = calc_nicehash_nonce(*piNonce, oWork.iResumeCnt); - else - pGpuCtx->Nonce = calc_start_nonce(oWork.iResumeCnt); + uint32_t h_per_round = pGpuCtx->rawIntensity; + size_t round_ctr = 0; assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID)); uint32_t target = oWork.iTarget32; @@ -201,6 +197,15 @@ void minethd::work_main() while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { + //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); + } + cl_uint results[0x100]; memset(results,0,sizeof(cl_uint)*(0x100)); -- cgit v1.1 From 98a4caa224f637d1e1025b5e385352e5156896a8 Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Sun, 15 Oct 2017 22:00:22 +0100 Subject: Implement changes suggested --- xmrstak/backend/amd/minethd.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'xmrstak/backend/amd') diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 048c3f0..8658e5c 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -198,7 +198,7 @@ void minethd::work_main() while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { //Allocate a new nonce every 16 rounds - if((++round_ctr & 0xF) == 0) + if((round_ctr++ & 0xF) == 0) { if(oWork.bNiceHash) pGpuCtx->Nonce = globalStates::inst().calc_start_nonce(pGpuCtx->Nonce & 0xFF000000u, h_per_round * 16); -- cgit v1.1 From d3c5911f8e486ac0d08c5160f6d9bd934c697b91 Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Tue, 17 Oct 2017 21:30:29 +0100 Subject: 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 +---- 3 files changed, 6 insertions(+), 7 deletions(-) (limited to 'xmrstak/backend/amd') 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]; -- cgit v1.1 From 265ef629a19128a2ae7faff31179a9f4816ce2a3 Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Tue, 17 Oct 2017 21:35:26 +0100 Subject: Fix GPU nicehash support --- xmrstak/backend/amd/minethd.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'xmrstak/backend/amd') diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index ca0b110..9d18860 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -170,7 +170,6 @@ void minethd::work_main() cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/); - pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39); globalStates::inst().iConsumeCnt++; while (bQuit == 0) @@ -195,6 +194,9 @@ void minethd::work_main() uint32_t target = oWork.iTarget32; XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target); + if(oWork.bNiceHash) + pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39); + while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { //Allocate a new nonce every 16 rounds -- cgit v1.1