[clang] 66d4dba - HIP: Directly call fmin/fmax builtins
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Thu Jun 8 10:50:13 PDT 2023
Author: Matt Arsenault
Date: 2023-06-08T13:49:40-04:00
New Revision: 66d4dbae02db2c0feb32ea844500f9758f078dd7
URL: https://github.com/llvm/llvm-project/commit/66d4dbae02db2c0feb32ea844500f9758f078dd7
DIFF: https://github.com/llvm/llvm-project/commit/66d4dbae02db2c0feb32ea844500f9758f078dd7.diff
LOG: HIP: Directly call fmin/fmax builtins
Added:
Modified:
clang/lib/Headers/__clang_hip_math.h
clang/test/Headers/__clang_hip_math.hip
clang/test/Headers/hip-header.hip
Removed:
################################################################################
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index 71cda10859ce2..7bfb98fba759f 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -247,10 +247,10 @@ float fmaf(float __x, float __y, float __z) {
}
__DEVICE__
-float fmaxf(float __x, float __y) { return __ocml_fmax_f32(__x, __y); }
+float fmaxf(float __x, float __y) { return __builtin_fmaxf(__x, __y); }
__DEVICE__
-float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); }
+float fminf(float __x, float __y) { return __builtin_fminf(__x, __y); }
__DEVICE__
float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }
@@ -796,10 +796,10 @@ double fma(double __x, double __y, double __z) {
}
__DEVICE__
-double fmax(double __x, double __y) { return __ocml_fmax_f64(__x, __y); }
+double fmax(double __x, double __y) { return __builtin_fmax(__x, __y); }
__DEVICE__
-double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); }
+double fmin(double __x, double __y) { return __builtin_fmin(__x, __y); }
__DEVICE__
double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }
@@ -1277,16 +1277,16 @@ __DEVICE__ int max(int __arg1, int __arg2) {
}
__DEVICE__
-float max(float __x, float __y) { return fmaxf(__x, __y); }
+float max(float __x, float __y) { return __builtin_fmaxf(__x, __y); }
__DEVICE__
-double max(double __x, double __y) { return fmax(__x, __y); }
+double max(double __x, double __y) { return __builtin_fmax(__x, __y); }
__DEVICE__
-float min(float __x, float __y) { return fminf(__x, __y); }
+float min(float __x, float __y) { return __builtin_fminf(__x, __y); }
__DEVICE__
-double min(double __x, double __y) { return fmin(__x, __y); }
+double min(double __x, double __y) { return __builtin_fmin(__x, __y); }
#if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
__host__ inline static int min(int __arg1, int __arg2) {
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index d4f87cd4f436a..29ccfd9c09b2e 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -979,13 +979,13 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: @test_fmaxf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmaxf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmax_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fmaxf(float x, float y) {
return fmaxf(x, y);
@@ -993,13 +993,13 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
// DEFAULT-LABEL: @test_fmax(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmax(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmax_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fmax(double x, double y) {
return fmax(x, y);
@@ -1007,13 +1007,13 @@ extern "C" __device__ double test_fmax(double x, double y) {
// DEFAULT-LABEL: @test_fminf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_fminf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmin_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fminf(float x, float y) {
return fminf(x, y);
@@ -1021,13 +1021,13 @@ extern "C" __device__ float test_fminf(float x, float y) {
// DEFAULT-LABEL: @test_fmin(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmin(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmin_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fmin(double x, double y) {
return fmin(x, y);
@@ -3694,13 +3694,13 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: @test_float_min(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_float_min(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmin_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_float_min(float x, float y) {
return min(x, y);
@@ -3708,13 +3708,13 @@ extern "C" __device__ float test_float_min(float x, float y) {
// DEFAULT-LABEL: @test_float_max(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_float_max(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmax_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_float_max(float x, float y) {
return max(x, y);
@@ -3722,13 +3722,13 @@ extern "C" __device__ float test_float_max(float x, float y) {
// DEFAULT-LABEL: @test_double_min(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_double_min(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmin_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_double_min(double x, double y) {
return min(x, y);
@@ -3736,13 +3736,13 @@ extern "C" __device__ double test_double_min(double x, double y) {
// DEFAULT-LABEL: @test_double_max(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_double_max(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmax_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_double_max(double x, double y) {
return max(x, y);
diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip
index 9aa7e1402e423..c11730f8717f5 100644
--- a/clang/test/Headers/hip-header.hip
+++ b/clang/test/Headers/hip-header.hip
@@ -123,7 +123,7 @@ __device__ float test_floor() {
}
// CHECK-LABEL: define{{.*}}@_Z8test_maxv
-// CHECK: call {{.*}}double @__ocml_fmax_f64(double {{.*}}, double
+// CHECK: call {{.*}}double @llvm.maxnum.f64(double {{.*}}, double
__device__ float test_max() {
return max(5, 6.0);
}
More information about the cfe-commits
mailing list