From 9652d69575f8cef33c54541835ddd9e2c220054a Mon Sep 17 00:00:00 2001 From: Aidan Belton-Schure Date: Thu, 10 Oct 2024 09:13:33 +0000 Subject: [PATCH] SWDEV-436099 Use new amdgcn_ballot builtin Change-Id: I024fabc6c5b3f39c66885eb7615953f4d0432e9a --- hipamd/include/hip/amd_detail/amd_warp_functions.h | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_warp_functions.h b/hipamd/include/hip/amd_detail/amd_warp_functions.h index de9b282b1c..a42440daa4 100644 --- a/hipamd/include/hip/amd_detail/amd_warp_functions.h +++ b/hipamd/include/hip/amd_detail/amd_warp_functions.h @@ -99,19 +99,16 @@ int __any(int predicate) { return __ockl_wfany_i32(predicate); } -// XXX from llvm/include/llvm/IR/InstrTypes.h -#define ICMP_NE 33 - __device__ inline unsigned long long int __ballot(int predicate) { - return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); + return __builtin_amdgcn_ballot_w64(predicate); } __device__ inline unsigned long long int __ballot64(int predicate) { - return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); + return __ballot(predicate); } // See amd_warp_sync_functions.h for an explanation of this preprocessor flag.