This is an archive of the discontinued LLVM Phabricator instance.

Differential D138395

HIP: Directly call fmin/fmax builtins
ClosedPublic

Authored by arsenm on Nov 20 2022, 8:48 PM.

Details

Reviewers
yaxunl
b-sumner
JonChesterfield

Diff Detail

Event Timeline

arsenm created this revision.Nov 20 2022, 8:48 PM
Herald added a project: Restricted Project. · View Herald TranscriptNov 20 2022, 8:48 PM
arsenm requested review of this revision.Nov 20 2022, 8:48 PM
Herald added a subscriber: wdng. · View Herald TranscriptNov 20 2022, 8:48 PM
arsenm added a parent revision: D138391: clang/HIP: Add new header test for math IR gen.Nov 20 2022, 8:48 PM
Harbormaster completed remote builds in B198716: Diff 476800.Nov 21 2022, 9:10 AM
arsenm added a comment.Jun 6 2023, 3:03 PM

ping

yaxunl accepted this revision.Jun 7 2023, 5:58 AM

LGTM. Thanks.

This revision is now accepted and ready to land.Jun 7 2023, 5:58 AM
arsenm closed this revision.Jun 8 2023, 10:50 AM

66d4dbae02db2c0feb32ea844500f9758f078dd7

Revision Contents

Diff 476800

clang/lib/Headers/__clang_hip_math.h

Show First 20 Lines • Show All 253 Lines • ▼ Show 20 Lines
float floorf(float __x) { return __ocml_floor_f32(__x); } float floorf(float __x) { return __ocml_floor_f32(__x); }
__DEVICE__ __DEVICE__
float fmaf(float __x, float __y, float __z) { float fmaf(float __x, float __y, float __z) {
return __builtin_fmaf(__x, __y, __z); return __builtin_fmaf(__x, __y, __z);
} }
__DEVICE__ __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__ __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__ __DEVICE__
float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); } float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }
__DEVICE__ __DEVICE__
float frexpf(float __x, int *__nptr) { float frexpf(float __x, int *__nptr) {
int __tmp; int __tmp;
#ifdef __OPENMP_AMDGCN__ #ifdef __OPENMP_AMDGCN__
▲ Show 20 Lines • Show All 536 Lines • ▼ Show 20 Lines
double floor(double __x) { return __ocml_floor_f64(__x); } double floor(double __x) { return __ocml_floor_f64(__x); }
__DEVICE__ __DEVICE__
double fma(double __x, double __y, double __z) { double fma(double __x, double __y, double __z) {
return __builtin_fma(__x, __y, __z); return __builtin_fma(__x, __y, __z);
} }
__DEVICE__ __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__ __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__ __DEVICE__
double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); } double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }
__DEVICE__ __DEVICE__
double frexp(double __x, int *__nptr) { double frexp(double __x, int *__nptr) {
int __tmp; int __tmp;
#ifdef __OPENMP_AMDGCN__ #ifdef __OPENMP_AMDGCN__
▲ Show 20 Lines • Show All 467 Lines • ▼ Show 20 Lines
__DEVICE__ int min(int __arg1, int __arg2) { __DEVICE__ int min(int __arg1, int __arg2) {
return (__arg1 < __arg2) ? __arg1 : __arg2; return (__arg1 < __arg2) ? __arg1 : __arg2;
} }
__DEVICE__ int max(int __arg1, int __arg2) { __DEVICE__ int max(int __arg1, int __arg2) {
return (__arg1 > __arg2) ? __arg1 : __arg2; return (__arg1 > __arg2) ? __arg1 : __arg2;
} }
__DEVICE__ __DEVICE__
float max(float __x, float __y) { return fmaxf(__x, __y); } float max(float __x, float __y) { return __builtin_fmaxf(__x, __y); }
__DEVICE__ __DEVICE__
double max(double __x, double __y) { return fmax(__x, __y); } double max(double __x, double __y) { return __builtin_fmax(__x, __y); }
__DEVICE__ __DEVICE__
float min(float __x, float __y) { return fminf(__x, __y); } float min(float __x, float __y) { return __builtin_fminf(__x, __y); }
__DEVICE__ __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__) #if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
__host__ inline static int min(int __arg1, int __arg2) { __host__ inline static int min(int __arg1, int __arg2) {
return std::min(__arg1, __arg2); return std::min(__arg1, __arg2);
} }
__host__ inline static int max(int __arg1, int __arg2) { __host__ inline static int max(int __arg1, int __arg2) {
return std::max(__arg1, __arg2); return std::max(__arg1, __arg2);
} }
#endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__) #endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
#endif #endif
#pragma pop_macro("__DEVICE__") #pragma pop_macro("__DEVICE__")
#pragma pop_macro("__RETURN_TYPE") #pragma pop_macro("__RETURN_TYPE")
#endif // __CLANG_HIP_MATH_H__ #endif // __CLANG_HIP_MATH_H__

clang/test/Headers/__clang_hip_math.hip

Show First 20 Lines • Show All 979 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret double [[TMP0]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test_fma_rn(double x, double y, double z) { extern "C" __device__ double test_fma_rn(double x, double y, double z) {
return __fma_rn(x, y, z); return __fma_rn(x, y, z);
} }
// DEFAULT-LABEL: @test_fmaxf( // DEFAULT-LABEL: @test_fmaxf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[CALL_I]] // DEFAULT-NEXT: ret float [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fmaxf( // FINITEONLY-LABEL: @test_fmaxf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[CALL_I]] // FINITEONLY-NEXT: ret float [[TMP0]]
// //
extern "C" __device__ float test_fmaxf(float x, float y) { extern "C" __device__ float test_fmaxf(float x, float y) {
return fmaxf(x, y); return fmaxf(x, y);
} }
// DEFAULT-LABEL: @test_fmax( // DEFAULT-LABEL: @test_fmax(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fmax( // FINITEONLY-LABEL: @test_fmax(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test_fmax(double x, double y) { extern "C" __device__ double test_fmax(double x, double y) {
return fmax(x, y); return fmax(x, y);
} }
// DEFAULT-LABEL: @test_fminf( // DEFAULT-LABEL: @test_fminf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[CALL_I]] // DEFAULT-NEXT: ret float [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fminf( // FINITEONLY-LABEL: @test_fminf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[CALL_I]] // FINITEONLY-NEXT: ret float [[TMP0]]
// //
extern "C" __device__ float test_fminf(float x, float y) { extern "C" __device__ float test_fminf(float x, float y) {
return fminf(x, y); return fminf(x, y);
} }
// DEFAULT-LABEL: @test_fmin( // DEFAULT-LABEL: @test_fmin(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fmin( // FINITEONLY-LABEL: @test_fmin(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test_fmin(double x, double y) { extern "C" __device__ double test_fmin(double x, double y) {
return fmin(x, y); return fmin(x, y);
} }
// DEFAULT-LABEL: @test_fmodf( // DEFAULT-LABEL: @test_fmodf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
▲ Show 20 Lines • Show All 2,540 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret double [[TMP0]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test__fma_rn(double x, double y, double z) { extern "C" __device__ double test__fma_rn(double x, double y, double z) {
return __fma_rn(x, y, z); return __fma_rn(x, y, z);
} }
// DEFAULT-LABEL: @test_float_min( // DEFAULT-LABEL: @test_float_min(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[CALL_I_I]] // DEFAULT-NEXT: ret float [[TMP0]]
// //
// FINITEONLY-LABEL: @test_float_min( // FINITEONLY-LABEL: @test_float_min(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[CALL_I_I]] // FINITEONLY-NEXT: ret float [[TMP0]]
// //
extern "C" __device__ float test_float_min(float x, float y) { extern "C" __device__ float test_float_min(float x, float y) {
return min(x, y); return min(x, y);
} }
// DEFAULT-LABEL: @test_float_max( // DEFAULT-LABEL: @test_float_max(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[CALL_I_I]] // DEFAULT-NEXT: ret float [[TMP0]]
// //
// FINITEONLY-LABEL: @test_float_max( // FINITEONLY-LABEL: @test_float_max(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[CALL_I_I]] // FINITEONLY-NEXT: ret float [[TMP0]]
// //
extern "C" __device__ float test_float_max(float x, float y) { extern "C" __device__ float test_float_max(float x, float y) {
return max(x, y); return max(x, y);
} }
// DEFAULT-LABEL: @test_double_min( // DEFAULT-LABEL: @test_double_min(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test_double_min( // FINITEONLY-LABEL: @test_double_min(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I_I]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test_double_min(double x, double y) { extern "C" __device__ double test_double_min(double x, double y) {
return min(x, y); return min(x, y);
} }
// DEFAULT-LABEL: @test_double_max( // DEFAULT-LABEL: @test_double_max(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test_double_max( // FINITEONLY-LABEL: @test_double_max(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I_I]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test_double_max(double x, double y) { extern "C" __device__ double test_double_max(double x, double y) {
return max(x, y); return max(x, y);
} }
// CHECK-LABEL: @test_int_min( // CHECK-LABEL: @test_int_min(
// CHECK-NEXT: entry: // CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.smin.i32(i32 [[X:%.*]], i32 [[Y:%.*]]) // CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.smin.i32(i32 [[X:%.*]], i32 [[Y:%.*]])
// CHECK-NEXT: ret i32 [[TMP0]] // CHECK-NEXT: ret i32 [[TMP0]]
Show All 13 Lines