From a8379c4d14c390d28411093ca0145ccbc99138dc Mon Sep 17 00:00:00 2001 From: Gerardo Hernandez Date: Mon, 12 Oct 2026 16:51:16 +0100 Subject: [PATCH 1/2] AIRUNTIME-2821 - fix two misuses of __builtin_amdgcn_is_invocable() --- .../hip_cooperative_groups_helper.h | 31 +++++++++++-------- 1 file changed, 18 insertions(+), 13 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index f654d8dc2406..bfda9fb56aaa 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -302,16 +302,19 @@ __CG_STATIC_QUALIFIER__ void sync() { // have i-th bit of x set and come before the current thread. __CG_STATIC_QUALIFIER__ unsigned int masked_bit_count(lane_mask x, unsigned int add = 0) { unsigned int counter = 0; - if (!__builtin_amdgcn_is_invocable(__builtin_amdgcn_mbcnt_lo)) __builtin_trap(); - if (static_cast(warpSize) == 32) { - counter = __builtin_amdgcn_mbcnt_lo(static_cast(x), add); + if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_mbcnt_lo)) { + if (static_cast(warpSize) == 32) { + counter = __builtin_amdgcn_mbcnt_lo(static_cast(x), add); + } else { + unsigned int lo = static_cast(x & 0xFFFFFFFF); + unsigned int hi = static_cast((x >> 32) & 0xFFFFFFFF); + counter = __builtin_amdgcn_mbcnt_lo(lo, add); + counter = __builtin_amdgcn_is_invocable(__builtin_amdgcn_mbcnt_hi) + ? __builtin_amdgcn_mbcnt_hi(hi, counter) + : counter; + } } else { - unsigned int lo = static_cast(x & 0xFFFFFFFF); - unsigned int hi = static_cast((x >> 32) & 0xFFFFFFFF); - counter = __builtin_amdgcn_mbcnt_lo(lo, add); - counter = __builtin_amdgcn_is_invocable(__builtin_amdgcn_mbcnt_hi) - ? __builtin_amdgcn_mbcnt_hi(hi, counter) - : counter; + __builtin_trap(); } return counter; @@ -437,10 +440,12 @@ template __CG_STATIC_QUALIFIER__ T* map_shared_rank(T* in, int rank __CG_STATIC_QUALIFIER__ unsigned int query_shared_rank(const void* in) { #if __has_builtin(__builtin_amdgcn_query_shared_rank) - if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_query_shared_rank)) - return static_cast( - __builtin_amdgcn_query_shared_rank((__attribute__((address_space(11))) const void*)in)); - __builtin_trap(); + if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_query_shared_rank)) { + return static_cast(__builtin_amdgcn_query_shared_rank((__attribute__((address_space(11))) const void*)in)); + } else { + __builtin_trap(); + return 0; + } #else return 0; #endif From cfced37d034929fde8dade924ea532fec42d3d16 Mon Sep 17 00:00:00 2001 From: Gerardo Hernandez Date: Tue, 13 Oct 2026 10:07:01 +0100 Subject: [PATCH 2/2] AIRUNTIME-2821 - remove extra whitespace --- .../include/hip/amd_detail/hip_cooperative_groups_helper.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index bfda9fb56aaa..94f50642f8e8 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -314,7 +314,7 @@ __CG_STATIC_QUALIFIER__ unsigned int masked_bit_count(lane_mask x, unsigned int : counter; } } else { - __builtin_trap(); + __builtin_trap(); } return counter;