Merge pull request #1062 from mhbliao/hliao/master/icmp

[hip] Re-implement ballot using AMDGCN builtins

hipamd/include/hip/hcc_detail/device_functions.h

@@ -736,13 +736,21 @@ int __any(int predicate) {
 __device__
 inline
 unsigned long long int __ballot(int predicate) {
+#if defined(__HCC__)
     return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE);
+#else
+     return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE);
+#endif
 }
 
 __device__
 inline
 unsigned long long int __ballot64(int predicate) {
+#if defined(__HCC__)
     return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE);
+#else
+     return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE);
+#endif
 }
 
 // hip.amdgcn.bc - lanemask