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

SWDEV-184380 Fix hcc compilation

Browse Source 
- `hcc` has no builtin. Need to invoke LLVM intrinsic directly.
This commit is contained in:
Michael LIAO
2019-03-26 15:11:42 -04:00
parent c20d233585
commit d355122bf9
1 changed files with 8 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/hcc_detail/device_functions.h
+8
View File
@@ -254,14 +254,22 @@ __device__ static inline float __hip_ds_permutef(int index, float src) {
template <int pattern>
__device__ static inline unsigned __hip_ds_swizzle_N(unsigned int src) {
union { int i; unsigned u; float f; } tmp; tmp.u = src;
#if defined(__HCC__)
tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern);
#else
tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern);
#endif
return tmp.u;
}
template <int pattern>
__device__ static inline float __hip_ds_swizzlef_N(float src) {
union { int i; unsigned u; float f; } tmp; tmp.f = src;
#if defined(__HCC__)
tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern);
#else
tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern);
#endif
return tmp.f;
}