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/
2 lines
8 lines
test/
Headers/
12 lines
40 lines

Diff 477238

clang/lib/Headers/__clang_hip_cmath.h

Show First 20 Lines • Show All 165 Lines • ▼ Show 20 Lines
// Notably missing above is nexttoward. We omit it because // Notably missing above is nexttoward. We omit it because
// ocml doesn't provide an implementation, and we don't want to be in the // ocml doesn't provide an implementation, and we don't want to be in the
// business of implementing tricky libm functions in this header. // business of implementing tricky libm functions in this header.
// Other functions. // Other functions.
__DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y, __DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y,
_Float16 __z) { _Float16 __z) {
return __ocml_fma_f16(__x, __y, __z); return __builtin_fmaf16(__x, __y, __z);
} }
__DEVICE__ __CONSTEXPR__ _Float16 pow(_Float16 __base, int __iexp) { __DEVICE__ __CONSTEXPR__ _Float16 pow(_Float16 __base, int __iexp) {
return __ocml_pown_f16(__base, __iexp); return __ocml_pown_f16(__base, __iexp);
} }
#ifndef __OPENMP_AMDGCN__ #ifndef __OPENMP_AMDGCN__
// BEGIN DEF_FUN and HIP_OVERLOAD // BEGIN DEF_FUN and HIP_OVERLOAD
▲ Show 20 Lines • Show All 660 Lines • Show Last 20 Lines

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 361 Lines • ▼ Show 20 Lines
} }
__DEVICE__ __DEVICE__
float __fmaf_rz(float __x, float __y, float __z) { float __fmaf_rz(float __x, float __y, float __z) {
return __ocml_fma_rtz_f32(__x, __y, __z); return __ocml_fma_rtz_f32(__x, __y, __z);
} }
#else #else
__DEVICE__ __DEVICE__
float __fmaf_rn(float __x, float __y, float __z) { float __fmaf_rn(float __x, float __y, float __z) {
return __ocml_fma_f32(__x, __y, __z); return __builtin_fmaf(__x, __y, __z);
} }
#endif #endif
#if defined OCML_BASIC_ROUNDED_OPERATIONS #if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__ __DEVICE__
float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); } float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); }
__DEVICE__ __DEVICE__
float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); } float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); }
▲ Show 20 Lines • Show All 161 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_cmath.hip

Show All 12 Lines
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/Inputs/include \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -ffinite-math-only -o - \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -ffinite-math-only -o - \
// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,FINITEONLY %s // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,FINITEONLY %s
// DEFAULT-LABEL: @test_fma_f16( // DEFAULT-LABEL: @test_fma_f16(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract half @__ocml_fma_f16(half noundef [[X:%.*]], half noundef [[Y:%.*]], half noundef [[Z:%.*]]) #[[ATTR6:[0-9]+]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
// DEFAULT-NEXT: ret half [[CALL_I]] // DEFAULT-NEXT: ret half [[TMP0]]
// //
// FINITEONLY-LABEL: @test_fma_f16( // FINITEONLY-LABEL: @test_fma_f16(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract half @__ocml_fma_f16(half noundef [[X:%.*]], half noundef [[Y:%.*]], half noundef [[Z:%.*]]) #[[ATTR6:[0-9]+]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
// FINITEONLY-NEXT: ret half [[CALL_I]] // FINITEONLY-NEXT: ret half [[TMP0]]
// //
extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y, extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
_Float16 z) { _Float16 z) {
return fma(x, y, z); return fma(x, y, z);
} }
// DEFAULT-LABEL: @test_pow_f16( // DEFAULT-LABEL: @test_pow_f16(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]] // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR6:[0-9]+]]
// DEFAULT-NEXT: ret half [[CALL_I]] // DEFAULT-NEXT: ret half [[CALL_I]]
// //
// FINITEONLY-LABEL: @test_pow_f16( // FINITEONLY-LABEL: @test_pow_f16(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR6:[0-9]+]]
// FINITEONLY-NEXT: ret half [[CALL_I]] // FINITEONLY-NEXT: ret half [[CALL_I]]
// //
extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) { extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) {
return pow(x, y); return pow(x, y);
} }

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,366 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret float [[DIV_I]] // FINITEONLY-NEXT: ret float [[DIV_I]]
// //
extern "C" __device__ float test___fdividef(float x, float y) { extern "C" __device__ float test___fdividef(float x, float y) {
return __fdividef(x, y); return __fdividef(x, y);
} }
// DEFAULT-LABEL: @test__fmaf_rn( // DEFAULT-LABEL: @test__fmaf_rn(
// 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_rn( // FINITEONLY-LABEL: @test__fmaf_rn(
// 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_rn(float x, float y, float z) { extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
return __fmaf_rn(x, y, z); return __fmaf_rn(x, y, z);
} }
// DEFAULT-LABEL: @test___fmul_rn( // DEFAULT-LABEL: @test___fmul_rn(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]] // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]]
▲ Show 20 Lines • Show All 241 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