diff options
| author | psychocrypt <psychocryptHPC@gmail.com> | 2018-04-27 21:45:40 +0200 |
|---|---|---|
| committer | Timothy Pearson <tpearson@raptorengineering.com> | 2018-06-04 21:07:11 +0000 |
| commit | 84471e3d48b0423410384d41c2ce467b19f2b73a (patch) | |
| tree | 6178551102088739478b093771b1fa92fbdc884e /xmrstak/backend/amd | |
| parent | 864baffb78f00773dd7e976a9c16aeef4af82e4d (diff) | |
| download | xmr-stak-84471e3d48b0423410384d41c2ce467b19f2b73a.zip xmr-stak-84471e3d48b0423410384d41c2ce467b19f2b73a.tar.gz | |
support stellite v4 fork
solve #1494
- add algorithm `cryptonight_v7_stellite` (internal named: `cryptonight_stellite`)
Diffstat (limited to 'xmrstak/backend/amd')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 2 | ||||
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 22 |
2 files changed, 16 insertions, 8 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 736b6ae..59f654f 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -1004,7 +1004,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar return(ERR_OCL_API); } - if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || cryptonight_stellite) { // Input if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS) diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 0738c04..9e2f03c 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -553,8 +553,8 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad, __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) , __global ulong *input #endif ) @@ -574,7 +574,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states } barrier(CLK_LOCAL_MEM_FENCE); -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) uint2 tweak1_2; #endif uint4 b_x; @@ -598,7 +599,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states b[1] = states[3] ^ states[7]; b_x = ((uint4 *)b)[0]; -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) tweak1_2 = as_uint2(input[4]); tweak1_2.s0 >>= 24; tweak1_2.s0 |= tweak1_2.s1 << 8; @@ -625,9 +627,15 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]); b_x ^= ((uint4 *)c)[0]; -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) uint table = 0x75310U; +// cryptonight_stellite +# if(ALGO == 7) + uint index = ((b_x.s2 >> 27) & 12) | ((b_x.s2 >> 23) & 2); +# else uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2); +# endif b_x.s2 ^= ((table >> index) & 0x30U) << 24; #endif Scratchpad[IDX((idx0 & MASK) >> 4)] = b_x; @@ -638,8 +646,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states a[1] += c[0] * as_ulong2(tmp).s0; a[0] += mul_hi(c[0], as_ulong2(tmp).s0); - -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) # if(ALGO == 6) uint2 ipbc_tmp = tweak1_2 ^ ((uint2 *)&(a[0]))[0]; |
