[Bf-blender-cvs] [390b9f13050] master: Fix Cycles HIP assuming warp size 32

[Brecht Van Lommel]

Commit: 390b9f1305059f5d8c7f944d44fc3e5821a3eb82
Author: Brecht Van Lommel
Date:   Wed Mar 16 15:03:38 2022 +0100
Branches: master
https://developer.blender.org/rB390b9f1305059f5d8c7f944d44fc3e5821a3eb82

Fix Cycles HIP assuming warp size 32

In HIP these masks are 64 bit, while in CUDA only 32 bit.

===================================================================

M	intern/cycles/kernel/device/hip/compat.h

===================================================================

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)