From adc4c3aabeb2b228760b7cf31d4b62299cb44291 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 21 Jun 2021 02:59:35 -0700 Subject: [PATCH] SWDEV-285388 - Bring back the inline asm till we have a proper compiler support Change-Id: Ibafe33dd8298f6e19f8b6a6d2aa220b633b76862 [ROCm/clr commit: f4680d98a97b8af49693335f9f7343e51c7c86ed] --- .../include/hip/amd_detail/amd_hip_atomic.h | 52 ++++++++++++++++--- .../include/hip/amd_detail/llvm_intrinsics.h | 19 ++----- .../hipamd/include/hip/amd_detail/math_fwd.h | 30 +++-------- 3 files changed, 54 insertions(+), 47 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h index 72c177d3dd..6a25533beb 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h @@ -281,16 +281,34 @@ __device__ inline unsigned int atomicInc(unsigned int* address, unsigned int val) { - return __builtin_amdgcn_atomic_inc32( - address, val, __ATOMIC_RELAXED, "agent"); + __device__ + extern + unsigned int __builtin_amdgcn_atomic_inc( + unsigned int*, + unsigned int, + unsigned int, + unsigned int, + bool) __asm("llvm.amdgcn.atomic.inc.i32.p0i32"); + + return __builtin_amdgcn_atomic_inc( + address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false); } __device__ inline unsigned int atomicDec(unsigned int* address, unsigned int val) { - return __builtin_amdgcn_atomic_dec32( - address, val, __ATOMIC_RELAXED, "agent"); + __device__ + extern + unsigned int __builtin_amdgcn_atomic_dec( + unsigned int*, + unsigned int, + unsigned int, + unsigned int, + bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32"); + + return __builtin_amdgcn_atomic_dec( + address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false); } __device__ @@ -580,16 +598,34 @@ __device__ inline unsigned int atomicInc(unsigned int* address, unsigned int val) { - return __builtin_amdgcn_atomic_inc32( - address, val, __ATOMIC_RELAXED, "agent"); + __device__ + extern + unsigned int __builtin_amdgcn_atomic_inc( + unsigned int*, + unsigned int, + unsigned int, + unsigned int, + bool) __asm("llvm.amdgcn.atomic.inc.i32.p0i32"); + + return __builtin_amdgcn_atomic_inc( + address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false); } __device__ inline unsigned int atomicDec(unsigned int* address, unsigned int val) { - return __builtin_amdgcn_atomic_dec32( - address, val, __ATOMIC_RELAXED, "agent"); + __device__ + extern + unsigned int __builtin_amdgcn_atomic_dec( + unsigned int*, + unsigned int, + unsigned int, + unsigned int, + bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32"); + + return __builtin_amdgcn_atomic_dec( + address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false); } __device__ diff --git a/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h b/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h index 5494b69f8b..c100b846c0 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h +++ b/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h @@ -22,7 +22,8 @@ THE SOFTWARE. /** * @file amd_detail/llvm_intrinsics.h - * @brief Contains declarations for wrapper functions for llvm intrinsics. + * @brief Contains declarations for wrapper functions for llvm intrinsics + * like llvm.amdgcn.s.barrier. */ #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_LLVM_INTRINSICS_H @@ -30,23 +31,11 @@ THE SOFTWARE. #include "hip/amd_detail/host_defines.h" -#if __has_builtin(__builtin_amdgcn_groupstaticsize) -__device__ -inline -unsigned __llvm_amdgcn_groupstaticsize() { - return __builtin_amdgcn_groupstaticsize(); -} -#else -// FIXME: This should be removed once the above builtin is available. +// FIXME: These should all be removed and proper builtins used. __device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); -#endif -template __device__ -inline -__T __llvm_amdgcn_ds_swizzle(__T index, const int pattern) { - return __builtin_amdgcn_ds_swizzle(index, pattern); -} +int __llvm_amdgcn_ds_swizzle(int index, int pattern) __asm("llvm.amdgcn.ds.swizzle"); #endif diff --git a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h index b299698c2f..ac46d537a8 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h +++ b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h @@ -388,29 +388,17 @@ __attribute__((const)) float __ocml_fma_rtz_f32(float, float, float); __device__ -inline __attribute__((const)) -float __llvm_amdgcn_cos_f32(float __x) { - return __builtin_amdgcn_cosf(__x); -} +float __llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32"); __device__ -inline __attribute__((const)) -float __llvm_amdgcn_rcp_f32(float __x) { - return __builtin_amdgcn_rcpf(__x); -} +float __llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32"); __device__ -inline __attribute__((const)) -float __llvm_amdgcn_rsq_f32(float __x) { - return __builtin_amdgcn_rsqf(__x); -} +float __llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32"); __device__ -inline __attribute__((const)) -float __llvm_amdgcn_sin_f32(float __x) { - return __builtin_amdgcn_sinf(__x); -} +float __llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32"); // END INTRINSICS // END FLOAT @@ -711,17 +699,11 @@ __attribute__((const)) double __ocml_fma_rtz_f64(double, double, double); __device__ -inline __attribute__((const)) -double __llvm_amdgcn_rcp_f64(double __x) { - return __builtin_amdgcn_rcp(__x); -} +double __llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64"); __device__ -inline __attribute__((const)) -double __llvm_amdgcn_rsq_f64(double __x) { - return __builtin_amdgcn_rsq(__x); -} +double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); // END INTRINSICS // END DOUBLE