diff options
| author | psychocrypt <psychocryptHPC@gmail.com> | 2018-04-22 21:11:54 +0200 |
|---|---|---|
| committer | Timothy Pearson <tpearson@raptorengineering.com> | 2018-06-04 21:07:11 +0000 |
| commit | f149b57458a05a9d5d67ad2834fa5e6c282fa76b (patch) | |
| tree | d6d9562f7789770593d85ef10a3551fe299794f7 /xmrstak/backend/amd/amd_gpu | |
| parent | 83db451571df19842d0791cfe11c98103543d84a (diff) | |
| download | xmr-stak-f149b57458a05a9d5d67ad2834fa5e6c282fa76b.zip xmr-stak-f149b57458a05a9d5d67ad2834fa5e6c282fa76b.tar.gz | |
add support for IPBC coin
- add algorithm `cryptonight_lite_v7_xor`
- update documentation
Diffstat (limited to 'xmrstak/backend/amd/amd_gpu')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 2 | ||||
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 21 |
2 files changed, 16 insertions, 7 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 03100d0..e44120b 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 ) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc) { // 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 d2ae1a7..0738c04 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 -#if(ALGO == 3 || ALGO == 5) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc +#if(ALGO == 3 || ALGO == 5 || ALGO == 6) , __global ulong *input #endif ) @@ -574,7 +574,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states } barrier(CLK_LOCAL_MEM_FENCE); -#if(ALGO == 3 || ALGO == 5) +#if(ALGO == 3 || ALGO == 5 || ALGO == 6) uint2 tweak1_2; #endif uint4 b_x; @@ -598,7 +598,7 @@ __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) +#if(ALGO == 3 || ALGO == 5 || ALGO == 6) tweak1_2 = as_uint2(input[4]); tweak1_2.s0 >>= 24; tweak1_2.s0 |= tweak1_2.s1 << 8; @@ -625,7 +625,7 @@ __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) +#if(ALGO == 3 || ALGO == 5 || ALGO == 6) uint table = 0x75310U; uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2); b_x.s2 ^= ((table >> index) & 0x30U) << 24; @@ -639,10 +639,19 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states a[0] += mul_hi(c[0], as_ulong2(tmp).s0); -#if(ALGO == 3 || ALGO == 5) +#if(ALGO == 3 || ALGO == 5 || ALGO == 6) + +# if(ALGO == 6) + uint2 ipbc_tmp = tweak1_2 ^ ((uint2 *)&(a[0]))[0]; + ((uint2 *)&(a[1]))[0] ^= ipbc_tmp; + Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0]; + ((uint2 *)&(a[1]))[0] ^= ipbc_tmp; +# else ((uint2 *)&(a[1]))[0] ^= tweak1_2; Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0]; ((uint2 *)&(a[1]))[0] ^= tweak1_2; +# endif + #else Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0]; #endif |
