SWDEV-1 - Use proper builtins for math intrinsics

Change-Id: If0e420332c718abe0e9f6316c467b6b29b80b183


[ROCm/clr commit: 14477c24d0]

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); }

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