From 390b9f1305059f5d8c7f944d44fc3e5821a3eb82 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Wed, 16 Mar 2022 15:03:38 +0100 Subject: Fix Cycles HIP assuming warp size 32 In HIP these masks are 64 bit, while in CUDA only 32 bit. --- intern/cycles/kernel/device/hip/compat.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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) -- cgit v1.2.3