From 22732aa90e8eb1219691875b4019257535cff462 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Wed, 16 May 2018 21:16:09 +0200 Subject: fix possible deadlock with Volta If CUDA 9.X is used and the miner is compiled for `sm_70` and used with Volta GPUs than the miner deadlocks if `threads` is not a multiple of `32`. - use `__activemask()` to get all active lanes --- xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 3d8af48..57b6ad0 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -187,7 +187,7 @@ __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_ unusedVar( ptr ); unusedVar( sub ); # if(__CUDACC_VER_MAJOR__ >= 9) - return __shfl_sync(0xFFFFFFFF, val, src, group_n ); + return __shfl_sync(__activemask(), val, src, group_n ); # else return __shfl( val, src, group_n ); # endif -- cgit v1.1