From d355122bf975dc25b368f1a84189b50ed8923004 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Tue, 26 Mar 2019 15:11:42 -0400 Subject: [PATCH] SWDEV-184380 Fix hcc compilation - `hcc` has no builtin. Need to invoke LLVM intrinsic directly. --- hipamd/include/hip/hcc_detail/device_functions.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index a24df22315..b88df2e168 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -254,14 +254,22 @@ __device__ static inline float __hip_ds_permutef(int index, float src) { template __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 __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; }