From 38e7eb0c74ba631c340d64ad190732603dfb16d6 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Mon, 1 Apr 2019 21:59:02 +0200 Subject: [PATCH] fix cn_r block 1802223 bug fix nvrtc deadlock with `cuda version != 10.1`: https://github.com/xmrig/xmrig-nvidia/issues/260 --- xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt b/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt index bcf495080..214114c7e 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt @@ -462,10 +462,10 @@ __global__ void CryptonightR_phase2( uint64_t bx0 = ((uint64_t*)(d_ctx_b + thread * 16))[sub]; uint64_t bx1 = ((uint64_t*)(d_ctx_b + thread * 16 + 4))[sub]; - uint32_t r0 = d_ctx_b[thread * 16 + 4 * 2]; - uint32_t r1 = d_ctx_b[thread * 16 + 4 * 2 + 1]; - uint32_t r2 = d_ctx_b[thread * 16 + 4 * 2 + 2]; - uint32_t r3 = d_ctx_b[thread * 16 + 4 * 2 + 3]; + volatile uint32_t r0 = d_ctx_b[thread * 16 + 4 * 2]; + volatile uint32_t r1 = d_ctx_b[thread * 16 + 4 * 2 + 1]; + volatile uint32_t r2 = d_ctx_b[thread * 16 + 4 * 2 + 2]; + volatile uint32_t r3 = d_ctx_b[thread * 16 + 4 * 2 + 3]; const int batchsize = (ITERATIONS * 2) >> ( 1 + bfactor ); const int start = partidx * batchsize;