SWDEV-436099 Use new amdgcn_ballot builtin

AidanBeltonS · AidanBeltonS · commit 9652d69575f8 · 2024-11-27T04:34:50.000-05:00
Change-Id: I024fabc6c5b3f39c66885eb7615953f4d0432e9a
diff --git 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.

Original file line number	Diff line number	Diff line change
`@@ -99,19 +99,16 @@ int __any(int predicate) {`
`99`	`99`	`return __ockl_wfany_i32(predicate);`
`100`	`100`	`}`
`101`	`101`
`102`		`-// XXX from llvm/include/llvm/IR/InstrTypes.h`
`103`		`-#define ICMP_NE 33`
`104`		`-`
`105`	`102`	`__device__`
`106`	`103`	`inline`
`107`	`104`	`unsigned long long int __ballot(int predicate) {`
`108`		`- return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE);`
	`105`	`+ return __builtin_amdgcn_ballot_w64(predicate);`
`109`	`106`	`}`
`110`	`107`
`111`	`108`	`__device__`
`112`	`109`	`inline`
`113`	`110`	`unsigned long long int __ballot64(int predicate) {`
`114`		`- return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE);`
	`111`	`+ return __ballot(predicate);`
`115`	`112`	`}`
`116`	`113`
`117`	`114`	`// See amd_warp_sync_functions.h for an explanation of this preprocessor flag.`