AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
[hip] Re-implement ballot using AMDGCN builtins

[ROCm/clr commit: 2eafa5dcf9]
This commit is contained in:
Maneesh Gupta
2019-05-03 17:48:19 +05:30
committed by GitHub
parent c73c864fdf e8de293fc5
commit 13b13c3493
1 changed files with 8 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
+8
View File
@@ -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