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 6a25533beb..72c177d3dd 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,34 +281,16 @@ __device__ inline unsigned int atomicInc(unsigned int* address, unsigned int val) { - __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); + return __builtin_amdgcn_atomic_inc32( + address, val, __ATOMIC_RELAXED, "agent"); } __device__ inline unsigned int atomicDec(unsigned int* address, unsigned int val) { - __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); + return __builtin_amdgcn_atomic_dec32( + address, val, __ATOMIC_RELAXED, "agent"); } __device__ @@ -598,34 +580,16 @@ __device__ inline unsigned int atomicInc(unsigned int* address, unsigned int val) { - __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); + return __builtin_amdgcn_atomic_inc32( + address, val, __ATOMIC_RELAXED, "agent"); } __device__ inline unsigned int atomicDec(unsigned int* address, unsigned int val) { - __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); + return __builtin_amdgcn_atomic_dec32( + address, val, __ATOMIC_RELAXED, "agent"); } __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 c100b846c0..8c626f7616 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h +++ b/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h @@ -22,8 +22,7 @@ THE SOFTWARE. /** * @file amd_detail/llvm_intrinsics.h - * @brief Contains declarations for wrapper functions for llvm intrinsics - * like llvm.amdgcn.s.barrier. + * @brief Contains declarations for wrapper functions for llvm intrinsics. */ #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_LLVM_INTRINSICS_H @@ -31,11 +30,17 @@ THE SOFTWARE. #include "hip/amd_detail/host_defines.h" -// FIXME: These should all be removed and proper builtins used. __device__ -unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); +inline +unsigned __llvm_amdgcn_groupstaticsize() { + return __builtin_amdgcn_groupstaticsize(); +} +template __device__ -int __llvm_amdgcn_ds_swizzle(int index, int pattern) __asm("llvm.amdgcn.ds.swizzle"); +inline +__T __llvm_amdgcn_ds_swizzle(__T index, const int pattern) { + return __builtin_amdgcn_ds_swizzle(index, pattern); +} #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 ac46d537a8..b299698c2f 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h +++ b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h @@ -388,17 +388,29 @@ __attribute__((const)) float __ocml_fma_rtz_f32(float, float, float); __device__ +inline __attribute__((const)) -float __llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32"); +float __llvm_amdgcn_cos_f32(float __x) { + return __builtin_amdgcn_cosf(__x); +} __device__ +inline __attribute__((const)) -float __llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32"); +float __llvm_amdgcn_rcp_f32(float __x) { + return __builtin_amdgcn_rcpf(__x); +} __device__ +inline __attribute__((const)) -float __llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32"); +float __llvm_amdgcn_rsq_f32(float __x) { + return __builtin_amdgcn_rsqf(__x); +} __device__ +inline __attribute__((const)) -float __llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32"); +float __llvm_amdgcn_sin_f32(float __x) { + return __builtin_amdgcn_sinf(__x); +} // END INTRINSICS // END FLOAT @@ -699,11 +711,17 @@ __attribute__((const)) double __ocml_fma_rtz_f64(double, double, double); __device__ +inline __attribute__((const)) -double __llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64"); +double __llvm_amdgcn_rcp_f64(double __x) { + return __builtin_amdgcn_rcp(__x); +} __device__ +inline __attribute__((const)) -double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); +double __llvm_amdgcn_rsq_f64(double __x) { + return __builtin_amdgcn_rsq(__x); +} // END INTRINSICS // END DOUBLE