github-jwinterm/xmr-stak
1
0
Fork 0
mirror of https://github.com/jwinterm/xmr-stak.git synced 2026-10-22 06:28:36 +02:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #68 from psychocrypt/topic-amdUse64bitTarget

Browse Source 
amd: use 64bit target
This commit is contained in:
fireice-uk
2017-10-27 15:07:48 +01:00
committed by GitHub
parent 712f7b7bdb 8888a60898
commit 847831e5f1
8 changed files with 22 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files
xmrstak/backend/amd/amd_gpu/gpu.cpp
+2 -2
View File
@@ -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;
xmrstak/backend/amd/amd_gpu/gpu.hpp
+1 -1
View File
@@ -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);
xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+17 -8
View File
@@ -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)
xmrstak/backend/amd/minethd.cpp
+1 -1
View File
@@ -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)
xmrstak/backend/miner_work.hpp
+1 -5
View File
@@ -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;
xmrstak/misc/executor.cpp
-2
View File
@@ -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;
xmrstak/net/jpsock.cpp
-2
View File
@@ -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)
{
xmrstak/net/msgstruct.hpp
-2
View File
@@ -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;