diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 9d86437b1..0b40c31c6 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -365,7 +365,7 @@ R"===( #if(STRIDED_INDEX==0) # define IDX(x) (x) #elif(STRIDED_INDEX==1) -# define IDX(x) ((x) * (Threads)) +# define IDX(x) (mul24(((uint)(x)), Threads)) #elif(STRIDED_INDEX==2) # define IDX(x) (((x) % MEM_CHUNK) + ((x) / MEM_CHUNK) * WORKSIZE * MEM_CHUNK) #elif(STRIDED_INDEX==3) @@ -867,7 +867,7 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states #pragma unroll 2 for(int i = 0, i1 = get_local_id(1); i < (MEMORY >> 7); ++i, i1 = (i1 + 16) % (MEMORY >> 4)) { - text ^= Scratchpad[IDX(i1)]; + text ^= Scratchpad[IDX((uint)i1)]; barrier(CLK_LOCAL_MEM_FENCE); text ^= *xin2_load; @@ -877,7 +877,7 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states *xin1_store = text; - text ^= Scratchpad[IDX(i1 + 8)]; + text ^= Scratchpad[IDX((uint)i1 + 8u)]; barrier(CLK_LOCAL_MEM_FENCE); text ^= *xin1_load; @@ -894,7 +894,7 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states #else #pragma unroll 2 for (int i = 0; i < (MEMORY >> 7); ++i) { - text ^= Scratchpad[IDX((i << 3) + get_local_id(1))]; + text ^= Scratchpad[IDX((uint)((i << 3) + get_local_id(1)))]; #pragma unroll 10 for(int j = 0; j < 10; ++j)