Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorBrecht Van Lommel <brecht@blender.org>2022-04-20 19:07:37 +0300
committerBrecht Van Lommel <brecht@blender.org>2022-04-20 19:09:23 +0300
commit8d2da45f98bc326d718b05172ff42300a4ab3988 (patch)
tree01f6a2a812068bbe83def5757673dc0020ebf9a8
parent3a6813ea651256cc2a056f82911feabdac554ea8 (diff)
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.
-rw-r--r--intern/cycles/kernel/device/hip/compat.h2
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 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)