[PATCH] D138504: clang/HIP: Remove __llvm_amdgcn_* wrapper hacks
Matt Arsenault via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Tue Nov 22 08:30:37 PST 2022
arsenm created this revision.
arsenm added reviewers: yaxunl, JonChesterfield, b-sumner.
Herald added a subscriber: kosarev.
Herald added a project: All.
arsenm requested review of this revision.
Herald added a subscriber: wdng.
These are leftover hacks from using asm declaratios to access
intrinsics.
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);
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D138504.477209.patch
Type: text/x-patch
Size: 2087 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20221122/fbecead9/attachment.bin>
More information about the cfe-commits
mailing list