From 8d2da45f98bc326d718b05172ff42300a4ab3988 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Wed, 20 Apr 2022 18:07:37 +0200 Subject: [PATCH] Revert "Fix Cycles HIP assuming warp size 32" This reverts commit 390b9f1305059f5d8c7f944d44fc3e5821a3eb82. It seems to break things on Linux for unknown reasons, so leave it out for now. A solution to this will be required for Vega cards though. --- 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 9c93d87fd87..667352ed12e 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) uint64_t(0xFFFFFFFFFFFFFFFF >> (64 - thread_warp)) +#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - 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)