From a60ade8c4ffb5bbfba08e797f75d04581264d486 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 28 Jan 2025 19:43:13 +0000 Subject: [PATCH] Explicitly select the correct `mbcnt` flavour, based on `warpSize`. --- hipamd/include/hip/amd_detail/amd_warp_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_warp_functions.h b/hipamd/include/hip/amd_detail/amd_warp_functions.h index a42440daa4..250874da9b 100644 --- a/hipamd/include/hip/amd_detail/amd_warp_functions.h +++ b/hipamd/include/hip/amd_detail/amd_warp_functions.h @@ -123,8 +123,8 @@ unsigned long long __activemask() { #endif // HIP_ENABLE_WARP_SYNC_BUILTINS __device__ static inline unsigned int __lane_id() { - return __builtin_amdgcn_mbcnt_hi( - -1, __builtin_amdgcn_mbcnt_lo(-1, 0)); + if (warpSize == 32) return __builtin_amdgcn_mbcnt_lo(-1, 0); + return __builtin_amdgcn_mbcnt_hi(-1, __builtin_amdgcn_mbcnt_lo(-1, 0)); } __device__