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.