diff options
Diffstat (limited to 'xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 22 |
1 files changed, 15 insertions, 7 deletions
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]; |
