This is an archive of the discontinued LLVM Phabricator instance.

Differential D138394

HIP: Directly call fma builtins
ClosedPublic

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

Details

Reviewers
yaxunl
JonChesterfield
b-sumner

Diff Detail

Event Timeline

arsenm created this revision.Nov 20 2022, 8:46 PM
Herald added a project: Restricted Project. · View Herald TranscriptNov 20 2022, 8:46 PM
arsenm requested review of this revision.Nov 20 2022, 8:46 PM
Herald added a subscriber: wdng. · View Herald TranscriptNov 20 2022, 8:46 PM
arsenm added a parent revision: D138391: clang/HIP: Add new header test for math IR gen.Nov 20 2022, 8:47 PM
Harbormaster completed remote builds in B198715: Diff 476799.Nov 21 2022, 9:09 AM
arsenm updated this revision to Diff 477213.Nov 22 2022, 8:37 AM

Missed one

Harbormaster completed remote builds in B198996: Diff 477213.Nov 22 2022, 8:37 AM
arsenm updated this revision to Diff 477238.Nov 22 2022, 9:52 AM

Catch another hidden in another header

Harbormaster completed remote builds in B199014: Diff 477238.Nov 22 2022, 9:53 AM
arsenm added a parent revision: D138509: clang/HIP: Add another math header test.Nov 22 2022, 9:54 AM
arsenm added a comment.Dec 22 2022, 5:10 AM

ping

yaxunl accepted this revision.Jan 10 2023, 7:26 AM
This revision is now accepted and ready to land.Jan 10 2023, 7:26 AM
arsenm closed this revision.Apr 26 2023, 4:27 PM

43fd46fda3c90b014e8a73c62f67af9543ea4d59

Revision Contents

PathSize
clang/
lib/
Headers/
6 lines
test/
Headers/
32 lines

Diff 476799

clang/lib/Headers/__clang_hip_math.h

Show First 20 Lines • Show All 249 Lines • ▼ Show 20 Lines
__DEVICE__ __DEVICE__
float fdividef(float __x, float __y) { return __x / __y; } float fdividef(float __x, float __y) { return __x / __y; }
__DEVICE__ __DEVICE__
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 __ocml_fma_f32(__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 __ocml_fmax_f32(__x, __y); }
__DEVICE__ __DEVICE__
float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); } float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); }
▲ Show 20 Lines • Show All 539 Lines • ▼ Show 20 Lines
__DEVICE__ __DEVICE__
double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); } double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); }
__DEVICE__ __DEVICE__
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 __ocml_fma_f64(__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 __ocml_fmax_f64(__x, __y); }
__DEVICE__ __DEVICE__
double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); } double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); }
▲ Show 20 Lines • Show All 442 Lines • ▼ Show 20 Lines
} }
__DEVICE__ __DEVICE__
double __fma_rz(double __x, double __y, double __z) { double __fma_rz(double __x, double __y, double __z) {
return __ocml_fma_rtz_f64(__x, __y, __z); return __ocml_fma_rtz_f64(__x, __y, __z);
} }
#else #else
__DEVICE__ __DEVICE__
double __fma_rn(double __x, double __y, double __z) { double __fma_rn(double __x, double __y, double __z) {
return __ocml_fma_f64(__x, __y, __z); return __builtin_fma(__x, __y, __z);
} }
#endif #endif
// END INTRINSICS // END INTRINSICS
// END DOUBLE // END DOUBLE
// C only macros // C only macros
#if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L #if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L
#define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x) #define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x)
▲ Show 20 Lines • Show All 49 Lines • Show Last 20 Lines

clang/test/Headers/__clang_hip_math.hip

Show First 20 Lines • Show All 937 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[CALL_I]]
// //
extern "C" __device__ double test_floor(double x) { extern "C" __device__ double test_floor(double x) {
return floor(x); return floor(x);
} }
// DEFAULT-LABEL: @test_fmaf( // DEFAULT-LABEL: @test_fmaf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
// DEFAULT-NEXT: ret float [[CALL_I]] // DEFAULT-NEXT: ret float [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fmaf( // FINITEONLY-LABEL: @test_fmaf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
// FINITEONLY-NEXT: ret float [[CALL_I]] // FINITEONLY-NEXT: ret float [[TMP0]]
// //
extern "C" __device__ float test_fmaf(float x, float y, float z) { extern "C" __device__ float test_fmaf(float x, float y, float z) {
return fmaf(x, y, z); return fmaf(x, y, z);
} }
// DEFAULT-LABEL: @test_fma( // DEFAULT-LABEL: @test_fma(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fma( // FINITEONLY-LABEL: @test_fma(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test_fma(double x, double y, double z) { extern "C" __device__ double test_fma(double x, double y, double z) {
return fma(x, y, z); return fma(x, y, z);
} }
// DEFAULT-LABEL: @test_fma_rn( // DEFAULT-LABEL: @test_fma_rn(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fma_rn( // FINITEONLY-LABEL: @test_fma_rn(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I]] // 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: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
▲ Show 20 Lines • Show All 2,582 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[CALL_I]]
// //
extern "C" __device__ double test___dsqrt_rn(double x) { extern "C" __device__ double test___dsqrt_rn(double x) {
return __dsqrt_rn(x); return __dsqrt_rn(x);
} }
// DEFAULT-LABEL: @test__fma_rn( // DEFAULT-LABEL: @test__fma_rn(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test__fma_rn( // FINITEONLY-LABEL: @test__fma_rn(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I]] // 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: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
▲ Show 20 Lines • Show All 69 Lines • Show Last 20 Lines