[PATCH] D138504: clang/HIP: Remove __llvm_amdgcn_* wrapper hacks

[Matt Arsenault via Phabricator via cfe-commits]

arsenm updated this revision to Diff 477219.
arsenm added a comment.

Remove some of the externally defined ones


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D138504/new/

https://reviews.llvm.org/D138504

Files:
  clang/lib/Headers/__clang_hip_libdevice_declares.h
  clang/lib/Headers/__clang_hip_math.h


Index: clang/lib/Headers/__clang_hip_math.h
===================================================================
--- clang/lib/Headers/__clang_hip_math.h
+++ clang/lib/Headers/__clang_hip_math.h
@@ -659,7 +659,7 @@
 #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__
Index: clang/lib/Headers/__clang_hip_libdevice_declares.h
===================================================================
--- clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -137,23 +137,6 @@
 __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
 
@@ -277,15 +260,6 @@
 __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);
@@ -305,7 +279,6 @@
 __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);
@@ -332,11 +305,6 @@
 __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);



Index: clang/lib/Headers/__clang_hip_math.h
===================================================================
--- clang/lib/Headers/__clang_hip_math.h
+++ clang/lib/Headers/__clang_hip_math.h
@@ -659,7 +659,7 @@
 #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__
Index: clang/lib/Headers/__clang_hip_libdevice_declares.h
===================================================================
--- clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -137,23 +137,6 @@
 __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
 
@@ -277,15 +260,6 @@
 __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);
@@ -305,7 +279,6 @@
 __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);
@@ -332,11 +305,6 @@
 __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);

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits