SWDEV-436099 Use new amdgcn_ballot builtin

Change-Id: I024fabc6c5b3f39c66885eb7615953f4d0432e9a


[ROCm/clr commit: 9652d69575]

projects/clr/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.