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 /intern/cycles/kernel/device/hip/compat.h
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.
Diffstat (limited to 'intern/cycles/kernel/device/hip/compat.h')
-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)