[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