diff options
| author | psychocrypt <psychocrypt@users.noreply.github.com> | 2017-12-01 22:37:35 +0100 |
|---|---|---|
| committer | psychocrypt <psychocrypt@users.noreply.github.com> | 2017-12-02 21:15:35 +0100 |
| commit | d10d20f1132a3c5d6dc51ed22a0381e460b579e8 (patch) | |
| tree | 12880be38994f39993bfa5e0d6b5c1b4c0312a53 /xmrstak/backend/amd/amd_gpu | |
| parent | 2920e9a3227da307b04ee23ecc5c63ecee4a224c (diff) | |
| download | xmr-stak-d10d20f1132a3c5d6dc51ed22a0381e460b579e8.zip xmr-stak-d10d20f1132a3c5d6dc51ed22a0381e460b579e8.tar.gz | |
add AMD scratchpad indexing option
Allow to change the indexing used to address the hash scratchpad memory.
- add option `strided_index` for each gpu
Diffstat (limited to 'xmrstak/backend/amd/amd_gpu')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 3 | ||||
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.hpp | 1 | ||||
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 19 |
3 files changed, 21 insertions, 2 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 879a2e4..42f6388 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -332,7 +332,8 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ char options[256]; snprintf(options, sizeof(options), - "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu", hasIterations, threadMemMask, int_port(ctx->workSize)); + "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d", + hasIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex ? 1 : 0); ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, options, NULL, NULL); if(ret != CL_SUCCESS) { diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index abbd08d..c17bac1 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -24,6 +24,7 @@ struct GpuContext size_t deviceIdx; size_t rawIntensity; size_t workSize; + int stridedIndex; /*Output vars*/ cl_device_id DeviceID; 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__) |
