From c77979a365afb493a79fec103f00776e1c1e9eec Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Fri, 23 Dec 2022 10:13:30 -0500
Subject: [PATCH] SWDEV-1 - Use proper builtins for math intrinsics

Change-Id: If0e420332c718abe0e9f6316c467b6b29b80b183


[ROCm/clr commit: 14477c24d0b9f7648e4b7e46e9baff8871d921db]
---
 .../hip/amd_detail/amd_math_functions.h       | 18 ++++++++---------
 .../hipamd/include/hip/amd_detail/math_fwd.h  | 20 -------------------
 2 files changed, 9 insertions(+), 29 deletions(-)

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