Skip to content

Commit

Permalink
HIP: Directly call fmin/fmax builtins
Browse files Browse the repository at this point in the history
  • Loading branch information
@arsenm
arsenm committed Jun 8, 2023
1 parent ab00b10 commit 66d4dba
Show file tree
Hide file tree
Showing 3 changed files with 41 additions and 41 deletions.
16 changes: 8 additions & 8 deletions clang/lib/Headers/__clang_hip_math.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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); }
Expand Down Expand Up @@ -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); }
Expand Down Expand Up @@ -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) {
Expand Down
64 changes: 32 additions & 32 deletions clang/test/Headers/__clang_hip_math.hip
View file
Original file line number Diff line number Diff line change
Expand Up @@ -979,55 +979,55 @@ 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);
}

// 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);
}

// 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);
}

// 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);
Expand Down Expand Up @@ -3694,55 +3694,55 @@ 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);
}

// 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);
}

// 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);
}

// 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);
Expand Down
2 changes: 1 addition & 1 deletion clang/test/Headers/hip-header.hip
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
Expand Down

0 comments on commit 66d4dba

Please sign in to comment.