diff options
| author | fireice-uk <fireice-uk@users.noreply.github.com> | 2017-12-04 15:47:31 +0000 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2017-12-04 15:47:31 +0000 |
| commit | a78b62d19756865879c54eac00c1c3f4d983d3a1 (patch) | |
| tree | dc5cd5b9f4bd1cb98e51e163b85e78cfad4edc67 /xmrstak/backend/amd/amd_gpu/opencl | |
| parent | 9c6bb11026de14c1a8d9eaecf5e70517d7a903e5 (diff) | |
| parent | cfe64725f063b63eafbeeef7e57cad7448cb3d15 (diff) | |
| download | xmr-stak-a78b62d19756865879c54eac00c1c3f4d983d3a1.zip xmr-stak-a78b62d19756865879c54eac00c1c3f4d983d3a1.tar.gz | |
Merge pull request #315 from psychocrypt/topic-amdScratchpadIndexing
add AMD scratchpad indexing option
Diffstat (limited to 'xmrstak/backend/amd/amd_gpu/opencl')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 19 |
1 files changed, 18 insertions, 1 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 966199b..255fcbb 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -411,7 +411,11 @@ void AESExpandKey256(uint *keybuf) } } -#define IDX(x) (x) +#if(STRIDED_INDEX==0) +# define IDX(x) (x) +#else +# define IDX(x) ((x) * (Threads)) +#endif __attribute__((reqd_work_group_size(WORKSIZE, 8, 1))) __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads) @@ -440,7 +444,12 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul if(gIdx < Threads) { states += 25 * gIdx; + +#if(STRIDED_INDEX==0) Scratchpad += gIdx * (ITERATIONS >> 2); +#else + Scratchpad += gIdx; +#endif ((ulong8 *)State)[0] = vload8(0, input); State[8] = input[8]; @@ -519,7 +528,11 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre if(gIdx < Threads) { states += 25 * gIdx; +#if(STRIDED_INDEX==0) Scratchpad += gIdx * (ITERATIONS >> 2); +#else + Scratchpad += gIdx; +#endif a[0] = states[0] ^ states[4]; b[0] = states[2] ^ states[6]; @@ -588,7 +601,11 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u if(gIdx < Threads) { states += 25 * gIdx; +#if(STRIDED_INDEX==0) Scratchpad += gIdx * (ITERATIONS >> 2); +#else + Scratchpad += gIdx; +#endif #if defined(__Tahiti__) || defined(__Pitcairn__) |
