diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_math_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_math_functions.h index 471f6ee362..3c17d298ea 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_math_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_math_functions.h @@ -640,22 +640,22 @@ inline float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); } __DEVICE__ inline -float __frcp_rd(float x) { return __llvm_amdgcn_rcp_f32(x); } +float __frcp_rd(float x) { return __builtin_amdgcn_rcpf(x); } #endif __DEVICE__ inline -float __frcp_rn(float x) { return __llvm_amdgcn_rcp_f32(x); } +float __frcp_rn(float x) { return __builtin_amdgcn_rcpf(x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ inline -float __frcp_ru(float x) { return __llvm_amdgcn_rcp_f32(x); } +float __frcp_ru(float x) { return __builtin_amdgcn_rcpf(x); } __DEVICE__ inline -float __frcp_rz(float x) { return __llvm_amdgcn_rcp_f32(x); } +float __frcp_rz(float x) { return __builtin_amdgcn_rcpf(x); } #endif __DEVICE__ inline -float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); } +float __frsqrt_rn(float x) { return __builtin_amdgcn_rsqf(x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ inline @@ -1155,18 +1155,18 @@ inline double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); } __DEVICE__ inline -double __drcp_rd(double x) { return __llvm_amdgcn_rcp_f64(x); } +double __drcp_rd(double x) { return __builtin_amdgcn_rcp(x); } #endif __DEVICE__ inline -double __drcp_rn(double x) { return __llvm_amdgcn_rcp_f64(x); } +double __drcp_rn(double x) { return __builtin_amdgcn_rcp(x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ inline -double __drcp_ru(double x) { return __llvm_amdgcn_rcp_f64(x); } +double __drcp_ru(double x) { return __builtin_amdgcn_rcp(x); } __DEVICE__ inline -double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); } +double __drcp_rz(double x) { return __builtin_amdgcn_rcp(x); } __DEVICE__ inline double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); } 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 050c88c114..9e999268ea 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h +++ b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h @@ -386,19 +386,6 @@ float __ocml_fma_rtp_f32(float, float, float); __device__ __attribute__((const)) float __ocml_fma_rtz_f32(float, float, float); - -__device__ -__attribute__((const)) -float __llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32"); -__device__ -__attribute__((const)) -float __llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32"); -__device__ -__attribute__((const)) -float __llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32"); -__device__ -__attribute__((const)) -float __llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32"); // END INTRINSICS // END FLOAT @@ -697,13 +684,6 @@ double __ocml_fma_rtp_f64(double, double, double); __device__ __attribute__((const)) double __ocml_fma_rtz_f64(double, double, double); - -__device__ -__attribute__((const)) -double __llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64"); -__device__ -__attribute__((const)) -double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); // END INTRINSICS // END DOUBLE