From 38e7eb0c74ba631c340d64ad190732603dfb16d6 Mon Sep 17 00:00:00 2001
From: psychocrypt <psychocryptHPC@gmail.com>
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 bcf4950..214114c 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;
--
GitLab