clang/HIP: Remove __llvm_amdgcn_* wrapper hacks

[arsenm]

Summary

These are leftover hacks from using asm declaratios to access
intrinsics.

[arsenm]

Remove some of the externally defined ones

[arsenm]

ping

[JonChesterfield]

Yep, thanks. Anything under __llvm is fair game to change and there's a find& replace fix for anything that decided to call these directly.

[arsenm]

f407a7399575a6821940973c54754d42e72dd9ce

clang/lib/Headers/__clang_hip_libdevice_declares.h

 __device__ __attribute__((const)) float __ocml_sqrt_rte_f32(float);
 __device__ __attribute__((const)) float __ocml_sqrt_rtn_f32(float);
 __device__ __attribute__((const)) float __ocml_sqrt_rtp_f32(float);
 __device__ __attribute__((const)) float __ocml_sqrt_rtz_f32(float);
 __device__ __attribute__((const)) float __ocml_fma_rte_f32(float, float, float);
 __device__ __attribute__((const)) float __ocml_fma_rtn_f32(float, float, float);
 __device__ __attribute__((const)) float __ocml_fma_rtp_f32(float, float, float);
 __device__ __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);
-}
-__device__ inline __attribute__((const)) float
-__llvm_amdgcn_rcp_f32(float __x) {
-  return __builtin_amdgcn_rcpf(__x);
-}
-__device__ inline __attribute__((const)) float
-__llvm_amdgcn_rsq_f32(float __x) {
-  return __builtin_amdgcn_rsqf(__x);
-}
-__device__ inline __attribute__((const)) float
-__llvm_amdgcn_sin_f32(float __x) {
-  return __builtin_amdgcn_sinf(__x);
-}
 // END INTRINSICS
 // END FLOAT
 
 // BEGIN DOUBLE
 __device__ __attribute__((const)) double __ocml_acos_f64(double);
 __device__ __attribute__((pure)) double __ocml_acosh_f64(double);
 __device__ __attribute__((const)) double __ocml_asin_f64(double);
 __device__ __attribute__((pure)) double __ocml_asinh_f64(double);
__device__ __attribute__((const)) double __ocml_fma_rte_f64(double, double,
                                                             double);
 __device__ __attribute__((const)) double __ocml_fma_rtn_f64(double, double,
                                                             double);
 __device__ __attribute__((const)) double __ocml_fma_rtp_f64(double, double,
                                                             double);
 __device__ __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);
-}
-__device__ inline __attribute__((const)) double
-__llvm_amdgcn_rsq_f64(double __x) {
-  return __builtin_amdgcn_rsq(__x);
-}
-
 __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
 __device__ _Float16 __ocml_cos_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_cvtrtn_f16_f32(float);
 __device__ __attribute__((const)) _Float16 __ocml_cvtrtp_f16_f32(float);
 __device__ __attribute__((const)) _Float16 __ocml_cvtrtz_f16_f32(float);
 __device__ __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_fma_f16(_Float16, _Float16,
                                                           _Float16);
 __device__ __attribute__((const)) _Float16 __ocml_fmax_f16(_Float16, _Float16);
 __device__ __attribute__((const)) _Float16 __ocml_fmin_f16(_Float16, _Float16);
 __device__ __attribute__((const)) _Float16 __ocml_fabs_f16(_Float16);
 __device__ __attribute__((const)) int __ocml_isinf_f16(_Float16);
 __device__ __attribute__((const)) int __ocml_isnan_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
-__device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
 __device__ _Float16 __ocml_sin_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_pown_f16(_Float16, int);
 
 typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));

 __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
 __device__ __attribute__((const))
 __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
 __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
 __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
-__device__ inline __2f16
-__llvm_amdgcn_rcp_2f16(__2f16 __x) // Not currently exposed by ROCDL.
-{
-  return (__2f16)(__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y));
-}
 __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
 __device__ __2f16 __ocml_sin_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
 
 #ifdef __cplusplus
 } // extern "C"
 #endif
 
 #endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__

clang/lib/Headers/__clang_hip_math.h

 __DEVICE__
 float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); }
 #else
 __DEVICE__
 float __frcp_rn(float __x) { return 1.0f / __x; }
 #endif
 
 __DEVICE__
-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__
 float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }
 __DEVICE__
 float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); }
 __DEVICE__
 float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }