diff options
| author | psychocrypt <psychocryptHPC@gmail.com> | 2018-05-16 21:16:09 +0200 |
|---|---|---|
| committer | psychocrypt <psychocryptHPC@gmail.com> | 2018-05-16 21:16:09 +0200 |
| commit | 4a4a2f856afe6581f661e234d04360909764a42f (patch) | |
| tree | 329915e4974d505f0f2c6a2cd92241c943319887 /xmrstak | |
| parent | 94d41580e0b0400e93a2f0226eb71476b891f4a7 (diff) | |
| download | xmr-stak-4a4a2f856afe6581f661e234d04360909764a42f.zip xmr-stak-4a4a2f856afe6581f661e234d04360909764a42f.tar.gz | |
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
Diffstat (limited to 'xmrstak')
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 2 |
1 files changed, 1 insertions, 1 deletions
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 |
