amd: use 64bit target

- remove 32bit target value
- use always 64bit target for amd backend

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;