[Bf-blender-cvs] [8d2da45f98b] master: Revert "Fix Cycles HIP assuming warp size 32"
Brecht Van Lommel
noreply at git.blender.org
Wed Apr 20 18:09:39 CEST 2022
Commit: 8d2da45f98bc326d718b05172ff42300a4ab3988
Author: Brecht Van Lommel
Date: Wed Apr 20 18:07:37 2022 +0200
Branches: master
https://developer.blender.org/rB8d2da45f98bc326d718b05172ff42300a4ab3988
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.
===================================================================
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 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)
More information about the Bf-blender-cvs
mailing list