From b0c2ac0a909d4ff4092df8a27050674d8d2e48e2 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 10 Mar 2017 08:40:59 -0600 Subject: [PATCH] Added architecture guards around __shfl, dpp and ds_permute device functions Change-Id: I10f9b08618fbf25b61c1932278fc5759e41c0d66 [ROCm/hip commit: 1567d20aa814cb8c440b356556de1c1188b37b44] --- .../hip/include/hip/hcc_detail/hip_runtime.h | 24 +++++++++++-------- 1 file changed, 14 insertions(+), 10 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 67c63103d3..6acc604909 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -226,6 +226,8 @@ __device__ int __all( int input); __device__ int __any( int input); __device__ unsigned long long int __ballot( int input); +#if __HIP_ARCH_GFX803__ == 1 + // warp shuffle functions #ifdef __cplusplus __device__ int __shfl(int input, int lane, int width=warpSize); @@ -247,6 +249,18 @@ __device__ float __shfl_down(float input, unsigned int lane_delta, int width); __device__ float __shfl_xor(float input, int lane_mask, int width); #endif +__device__ unsigned __hip_ds_bpermute(int index, unsigned src); +__device__ float __hip_ds_bpermutef(int index, float src); +__device__ unsigned __hip_ds_permute(int index, unsigned src); +__device__ float __hip_ds_permutef(int index, float src); + +__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern); +__device__ float __hip_ds_swizzlef(float src, int pattern); + +__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl); + +#endif + __host__ __device__ int min(int arg1, int arg2); __host__ __device__ int max(int arg1, int arg2); @@ -321,16 +335,6 @@ __device__ static inline void __threadfence(void) { //__device__ void __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details"))); __device__ void __threadfence_system(void) ; -__device__ unsigned __hip_ds_bpermute(int index, unsigned src); -__device__ float __hip_ds_bpermutef(int index, float src); -__device__ unsigned __hip_ds_permute(int index, unsigned src); -__device__ float __hip_ds_permutef(int index, float src); - -__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern); -__device__ float __hip_ds_swizzlef(float src, int pattern); - -__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl); - // doxygen end Fence Fence /** * @}