diff options
| author | psychocrypt <psychocryptHPC@gmail.com> | 2018-05-16 21:16:09 +0200 |
|---|---|---|
| committer | Timothy Pearson <tpearson@raptorengineering.com> | 2018-06-04 21:07:11 +0000 |
| commit | 22732aa90e8eb1219691875b4019257535cff462 (patch) | |
| tree | bf9c7fc06c4ec638f1a6f2130b7d836cef7a1e8d | |
| parent | a50ac4590ebc7aca52ebff686dd774b5ac0b3af3 (diff) | |
| download | xmr-stak-22732aa90e8eb1219691875b4019257535cff462.zip xmr-stak-22732aa90e8eb1219691875b4019257535cff462.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
| -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 |
