blender
/
blender
136
404
Fork 586
Code Issues 6.2k Pull Requests 752 Projects 19 Wiki Activity

Fix Cycles HIP assuming warp size 32 

In HIP these masks are 64 bit, while in CUDA only 32 bit.
This commit is contained in:
Brecht Van Lommel 2022-03-16 15:03:38 +01:00
parent 076079454f
commit 390b9f1305
Notes: blender-bot 2023-02-14 10:21:15 +01:00 
Referenced by commit 8d2da45f98, Revert "Fix Cycles HIP assuming warp size 32"
1 changed files with 1 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
intern/cycles/kernel/device/hip/compat.h
View File

@ -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) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
#define ccl_gpu_thread_mask(thread_warp) uint64_t(0xFFFFFFFFFFFFFFFF >> (64 - 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)