[clang] 43fd46f - HIP: Directly call fma builtins
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 26 16:27:19 PDT 2023
Author: Matt Arsenault
Date: 2023-04-26T19:27:11-04:00
New Revision: 43fd46fda3c90b014e8a73c62f67af9543ea4d59
URL: https://github.com/llvm/llvm-project/commit/43fd46fda3c90b014e8a73c62f67af9543ea4d59
DIFF: https://github.com/llvm/llvm-project/commit/43fd46fda3c90b014e8a73c62f67af9543ea4d59.diff
LOG: HIP: Directly call fma builtins
Added:
Modified:
clang/lib/Headers/__clang_hip_math.h
clang/test/Headers/__clang_hip_math.hip
Removed:
################################################################################
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index 537dd0fca8702..a4e557e45e5a0 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -243,7 +243,7 @@ float floorf(float __x) { return __ocml_floor_f32(__x); }
__DEVICE__
float fmaf(float __x, float __y, float __z) {
- return __ocml_fma_f32(__x, __y, __z);
+ return __builtin_fmaf(__x, __y, __z);
}
__DEVICE__
@@ -621,7 +621,7 @@ float __fmaf_rz(float __x, float __y, float __z) {
#else
__DEVICE__
float __fmaf_rn(float __x, float __y, float __z) {
- return __ocml_fma_f32(__x, __y, __z);
+ return __builtin_fmaf(__x, __y, __z);
}
#endif
@@ -799,7 +799,7 @@ double floor(double __x) { return __ocml_floor_f64(__x); }
__DEVICE__
double fma(double __x, double __y, double __z) {
- return __ocml_fma_f64(__x, __y, __z);
+ return __builtin_fma(__x, __y, __z);
}
__DEVICE__
@@ -1258,7 +1258,7 @@ double __fma_rz(double __x, double __y, double __z) {
#else
__DEVICE__
double __fma_rn(double __x, double __y, double __z) {
- return __ocml_fma_f64(__x, __y, __z);
+ return __builtin_fma(__x, __y, __z);
}
#endif
// END INTRINSICS
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index ec774c43f8124..e546ce5ac24ab 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -767,6 +767,7 @@ extern "C" __device__ float test_exp2f(float x) {
return exp2f(x);
}
+//
// DEFAULT-LABEL: @test_exp2(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -937,13 +938,13 @@ extern "C" __device__ double test_floor(double x) {
// DEFAULT-LABEL: @test_fmaf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmaf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fma_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fmaf(float x, float y, float z) {
return fmaf(x, y, z);
@@ -951,13 +952,13 @@ extern "C" __device__ float test_fmaf(float x, float y, float z) {
// DEFAULT-LABEL: @test_fma(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fma(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fma_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fma(double x, double y, double z) {
return fma(x, y, z);
@@ -965,13 +966,13 @@ extern "C" __device__ double test_fma(double x, double y, double z) {
// DEFAULT-LABEL: @test_fma_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fma_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fma_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fma_rn(double x, double y, double z) {
return __fma_rn(x, y, z);
@@ -3414,13 +3415,13 @@ extern "C" __device__ float test___fdividef(float x, float y) {
// DEFAULT-LABEL: @test__fmaf_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test__fmaf_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fma_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
return __fmaf_rn(x, y, z);
@@ -3692,13 +3693,13 @@ extern "C" __device__ double test___dsqrt_rn(double x) {
// DEFAULT-LABEL: @test__fma_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test__fma_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fma_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test__fma_rn(double x, double y, double z) {
return __fma_rn(x, y, z);
More information about the cfe-commits
mailing list