diff options
| author | Brecht Van Lommel <brecht@blender.org> | 2022-03-16 17:03:38 +0300 |
|---|---|---|
| committer | Brecht Van Lommel <brecht@blender.org> | 2022-03-16 20:05:48 +0300 |
| commit | 390b9f1305059f5d8c7f944d44fc3e5821a3eb82 (patch) | |
| tree | db3293fa4d6e0a91ccc905e9115140c8be3c022a /intern/cycles/kernel/device/hip/compat.h | |
| parent | 076079454fdf4e6768d0c639ab5dbfc16fca0d50 (diff) | |
Fix Cycles HIP assuming warp size 32
In HIP these masks are 64 bit, while in CUDA only 32 bit.
Diffstat (limited to 'intern/cycles/kernel/device/hip/compat.h')
| -rw-r--r-- | intern/cycles/kernel/device/hip/compat.h | 2 |
1 files changed, 1 insertions, 1 deletions
diff --git a/intern/cycles/kernel/device/hip/compat.h b/intern/cycles/kernel/device/hip/compat.h index 94d9d1273e8..29fbc119cd1 100644 --- a/intern/cycles/kernel/device/hip/compat.h +++ b/intern/cycles/kernel/device/hip/compat.h @@ -62,7 +62,7 @@ typedef unsigned long long uint64_t; #define ccl_gpu_block_idx_x (blockIdx.x) #define ccl_gpu_grid_dim_x (gridDim.x) #define ccl_gpu_warp_size (warpSize) -#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp)) +#define ccl_gpu_thread_mask(thread_warp) uint64_t(0xFFFFFFFFFFFFFFFF >> (64 - thread_warp)) #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x) #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x) |