From ba5a9a5395bb35399ba2e44a72cc0134bda58b09 Mon Sep 17 00:00:00 2001 From: "Hernandez, Gerardo" Date: Tue, 22 Apr 2025 15:17:58 +0100 Subject: [PATCH] SWDEV-420237 - Fix reduce sync operations when masks are divergent (#181) Do not use __ockl_activelane_u32() to calculate the index of the lane within the mask, as that would not work with divergent masks that have other bits on before the associated lane. [ROCm/clr commit: 1a8d766836d0731d936aa9129aa524424363da20] --- .../hipamd/include/hip/amd_detail/amd_warp_sync_functions.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h index 9ba74fd3e1..a47ad1ea9d 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h @@ -373,6 +373,7 @@ __device__ inline T __reduce_op_sync(MaskT mask, T val, BinaryOp op, WfReduce wf // the number of bits in the mask is a power of 2 numIterations -= 1; + maskIdx = __popcll(((1ul << laneId) - 1) & mask); mask >>= laneId; mask >>= 1ul; @@ -381,8 +382,6 @@ __device__ inline T __reduce_op_sync(MaskT mask, T val, BinaryOp op, WfReduce wf else __builtin_memcpy(&result, &val, sizeof(T)); - maskIdx = __ockl_activelane_u32(); - // add the values from the lanes using a reduction tree (first the threads with even-numbered // lanes, then multiples of 4, then 8, ... while (numIterations) {