From 8888a608987c85870a99bfa99bd6980018288666 Mon Sep 17 00:00:00 2001
From: psychocrypt <psychocrypt@users.noreply.github.com>
Date: Tue, 24 Oct 2017 23:22:28 +0200
Subject: [PATCH] amd: use 64bit target
- remove 32bit target value
- use always 64bit target for amd backend
---
xmrstak/backend/amd/amd_gpu/gpu.cpp | 4 +--
xmrstak/backend/amd/amd_gpu/gpu.hpp | 2 +-
.../backend/amd/amd_gpu/opencl/cryptonight.cl | 25 +++++++++++++------
xmrstak/backend/amd/minethd.cpp | 2 +-
xmrstak/backend/miner_work.hpp | 6 +----
xmrstak/misc/executor.cpp | 2 --
xmrstak/net/jpsock.cpp | 2 --
xmrstak/net/msgstruct.hpp | 2 --
8 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 37adc5b..3575854 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -645,7 +645,7 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx)
return ERR_SUCCESS;
}
-size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint32_t target)
+size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target)
{
cl_int ret;
@@ -787,7 +787,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint32_t tar
}
// Target
- if((ret = clSetKernelArg(ctx->Kernels[i + 3], 3, sizeof(cl_uint), &target)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[i + 3], 3, sizeof(cl_ulong), &target)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 3);
return ERR_OCL_API;
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp
index c2d708d..123de01 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.hpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp
@@ -45,7 +45,7 @@ int getAMDPlatformIdx();
std::vector<GpuContext> getAMDDevices(int index);
size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx);
-size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint32_t target);
+size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target);
size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput);
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
index a1642c4..a6a5910 100644
--- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
@@ -665,7 +665,7 @@ R"===(
#define VSWAP4(x) ((((x) >> 24) & 0xFFU) | (((x) >> 8) & 0xFF00U) | (((x) << 8) & 0xFF0000U) | (((x) << 24) & 0xFF000000U))
-__kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global uint *output, uint Target, ulong Threads)
+__kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global uint *output, ulong Target, ulong Threads)
{
const ulong idx = get_global_id(0) - get_global_offset(0);
@@ -713,7 +713,9 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u
//vstore8(p, 0, output);
- if(as_uint16(p).s7 <= Target)
+ // Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
+ // and expect an accurate result for target > 32-bit without implementing carries
+ if(p.s3 <= Target)
{
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
@@ -725,7 +727,7 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u
#define SWAP8(x) as_ulong(as_uchar8(x).s76543210)
-__kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint *output, uint Target, ulong Threads)
+__kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint *output, ulong Target, ulong Threads)
{
const uint idx = get_global_id(0) - get_global_offset(0);
@@ -785,7 +787,9 @@ __kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint
//output[2] = h7h;
//output[3] = h7l;
- if(as_uint2(h7l).s1 <= Target)
+ // Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
+ // and expect an accurate result for target > 32-bit without implementing carries
+ if(h7l <= Target)
{
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
@@ -796,7 +800,7 @@ __kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint
#define SWAP4(x) as_uint(as_uchar4(x).s3210)
-__kernel void Blake(__global ulong *states, __global uint *BranchBuf, __global uint *output, uint Target, ulong Threads)
+__kernel void Blake(__global ulong *states, __global uint *BranchBuf, __global uint *output, ulong Target, ulong Threads)
{
const uint idx = get_global_id(0) - get_global_offset(0);
@@ -859,7 +863,10 @@ __kernel void Blake(__global ulong *states, __global uint *BranchBuf, __global u
for(int i = 0; i < 8; ++i) h[i] = SWAP4(h[i]);
- if(h[7] <= Target)
+ // Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
+ // and expect an accurate result for target > 32-bit without implementing carries
+ uint2 t = (uint2)(h[6],h[7]);
+ if( as_ulong(t) <= Target)
{
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
@@ -868,7 +875,7 @@ __kernel void Blake(__global ulong *states, __global uint *BranchBuf, __global u
}
}
-__kernel void Groestl(__global ulong *states, __global uint *BranchBuf, __global uint *output, uint Target, ulong Threads)
+__kernel void Groestl(__global ulong *states, __global uint *BranchBuf, __global uint *output, ulong Target, ulong Threads)
{
const uint idx = get_global_id(0) - get_global_offset(0);
@@ -917,7 +924,9 @@ __kernel void Groestl(__global ulong *states, __global uint *BranchBuf, __global
for(int i = 0; i < 8; ++i) State[i] ^= tmp[i];
- if(as_uint2(State[7]).s1 <= Target)
+ // Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
+ // and expect an accurate result for target > 32-bit without implementing carries
+ if(State[7] <= Target)
{
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index 9d18860..4cbac6d 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -191,7 +191,7 @@ void minethd::work_main()
size_t round_ctr = 0;
assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID));
- uint32_t target = oWork.iTarget32;
+ uint64_t target = oWork.iTarget;
XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target);
if(oWork.bNiceHash)
diff --git a/xmrstak/backend/miner_work.hpp b/xmrstak/backend/miner_work.hpp
index 6b5720c..4bfe429 100644
--- a/xmrstak/backend/miner_work.hpp
+++ b/xmrstak/backend/miner_work.hpp
@@ -16,8 +16,6 @@ namespace xmrstak
uint8_t bWorkBlob[112];
uint32_t iWorkSize;
uint64_t iTarget;
- // \todo remove workaround needed for amd
- uint32_t iTarget32;
bool bNiceHash;
bool bStall;
size_t iPoolId;
@@ -41,7 +39,6 @@ namespace xmrstak
iWorkSize = from.iWorkSize;
iTarget = from.iTarget;
- iTarget32 = from.iTarget32;
bNiceHash = from.bNiceHash;
bStall = from.bStall;
iPoolId = from.iPoolId;
@@ -53,7 +50,7 @@ namespace xmrstak
return *this;
}
- miner_work(miner_work&& from) : iWorkSize(from.iWorkSize), iTarget(from.iTarget),iTarget32(from.iTarget32),
+ miner_work(miner_work&& from) : iWorkSize(from.iWorkSize), iTarget(from.iTarget),
bStall(from.bStall), iPoolId(from.iPoolId)
{
assert(iWorkSize <= sizeof(bWorkBlob));
@@ -67,7 +64,6 @@ namespace xmrstak
iWorkSize = from.iWorkSize;
iTarget = from.iTarget;
- iTarget32 = from.iTarget32;
bNiceHash = from.bNiceHash;
bStall = from.bStall;
iPoolId = from.iPoolId;
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index c518e19..7fc46e4 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -238,8 +238,6 @@ void executor::on_pool_have_job(size_t pool_id, pool_job& oPoolJob)
xmrstak::miner_work oWork(oPoolJob.sJobID, oPoolJob.bWorkBlob, oPoolJob.iWorkLen, oPoolJob.iTarget,
pool_id != dev_pool_id && ::jconf::inst()->NiceHashMode(), pool_id);
-
- oWork.iTarget32 = oPoolJob.iTarget32;
xmrstak::pool_data dat;
dat.iSavedNonce = oPoolJob.iSavedNonce;
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index 7735921..68b495d 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -396,8 +396,6 @@ bool jpsock::process_pool_job(const opq_json_val* params)
oPoolJob.iTarget = t32_to_t64(iTempInt);
- oPoolJob.iTarget32 = iTempInt;
-
}
else if(target_slen <= 16)
{
diff --git a/xmrstak/net/msgstruct.hpp b/xmrstak/net/msgstruct.hpp
index 8e8254b..82b59c1 100644
--- a/xmrstak/net/msgstruct.hpp
+++ b/xmrstak/net/msgstruct.hpp
@@ -12,8 +12,6 @@ struct pool_job
char sJobID[64];
uint8_t bWorkBlob[112];
uint64_t iTarget;
- // \todo remove workaround needed for amd
- uint32_t iTarget32;
uint32_t iWorkLen;
uint32_t iSavedNonce;
--
GitLab