SWDEV-184380 Fix hcc compilation

- `hcc` has no builtin. Need to invoke LLVM intrinsic directly.
这个提交包含在:
Michael LIAO 于 2019-03-26 15:20:17 -04:00 提交
1 父节点 30b5c02ec4
当前提交 939d153c1a
1 个文件已更改
+8
@@ -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;
}