Author: Matt Arsenault Date: 2023-04-26T19:27:11-04:00 New Revision: 43fd46fda3c90b014e8a73c62f67af9543ea4d59
URL: https://github.com/llvm/llvm-project/commit/43fd46fda3c90b014e8a73c62f67af9543ea4d59 DIFF: https://github.com/llvm/llvm-project/commit/43fd46fda3c90b014e8a73c62f67af9543ea4d59.diff LOG: HIP: Directly call fma builtins Added: Modified: clang/lib/Headers/__clang_hip_math.h clang/test/Headers/__clang_hip_math.hip Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index 537dd0fca8702..a4e557e45e5a0 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -243,7 +243,7 @@ float floorf(float __x) { return __ocml_floor_f32(__x); } __DEVICE__ float fmaf(float __x, float __y, float __z) { - return __ocml_fma_f32(__x, __y, __z); + return __builtin_fmaf(__x, __y, __z); } __DEVICE__ @@ -621,7 +621,7 @@ float __fmaf_rz(float __x, float __y, float __z) { #else __DEVICE__ float __fmaf_rn(float __x, float __y, float __z) { - return __ocml_fma_f32(__x, __y, __z); + return __builtin_fmaf(__x, __y, __z); } #endif @@ -799,7 +799,7 @@ double floor(double __x) { return __ocml_floor_f64(__x); } __DEVICE__ double fma(double __x, double __y, double __z) { - return __ocml_fma_f64(__x, __y, __z); + return __builtin_fma(__x, __y, __z); } __DEVICE__ @@ -1258,7 +1258,7 @@ double __fma_rz(double __x, double __y, double __z) { #else __DEVICE__ double __fma_rn(double __x, double __y, double __z) { - return __ocml_fma_f64(__x, __y, __z); + return __builtin_fma(__x, __y, __z); } #endif // END INTRINSICS diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index ec774c43f8124..e546ce5ac24ab 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -767,6 +767,7 @@ extern "C" __device__ float test_exp2f(float x) { return exp2f(x); } +// // DEFAULT-LABEL: @test_exp2( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]] @@ -937,13 +938,13 @@ extern "C" __device__ double test_floor(double x) { // DEFAULT-LABEL: @test_fmaf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]] -// DEFAULT-NEXT: ret float [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]]) +// DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: @test_fmaf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fma_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] -// FINITEONLY-NEXT: ret float [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]]) +// FINITEONLY-NEXT: ret float [[TMP0]] // extern "C" __device__ float test_fmaf(float x, float y, float z) { return fmaf(x, y, z); @@ -951,13 +952,13 @@ extern "C" __device__ float test_fmaf(float x, float y, float z) { // DEFAULT-LABEL: @test_fma( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]] -// DEFAULT-NEXT: ret double [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// DEFAULT-NEXT: ret double [[TMP0]] // // FINITEONLY-LABEL: @test_fma( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fma_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// FINITEONLY-NEXT: ret double [[TMP0]] // extern "C" __device__ double test_fma(double x, double y, double z) { return fma(x, y, z); @@ -965,13 +966,13 @@ extern "C" __device__ double test_fma(double x, double y, double z) { // DEFAULT-LABEL: @test_fma_rn( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]] -// DEFAULT-NEXT: ret double [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// DEFAULT-NEXT: ret double [[TMP0]] // // FINITEONLY-LABEL: @test_fma_rn( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fma_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// FINITEONLY-NEXT: ret double [[TMP0]] // extern "C" __device__ double test_fma_rn(double x, double y, double z) { return __fma_rn(x, y, z); @@ -3414,13 +3415,13 @@ extern "C" __device__ float test___fdividef(float x, float y) { // DEFAULT-LABEL: @test__fmaf_rn( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]] -// DEFAULT-NEXT: ret float [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]]) +// DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: @test__fmaf_rn( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fma_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] -// FINITEONLY-NEXT: ret float [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]]) +// FINITEONLY-NEXT: ret float [[TMP0]] // extern "C" __device__ float test__fmaf_rn(float x, float y, float z) { return __fmaf_rn(x, y, z); @@ -3692,13 +3693,13 @@ extern "C" __device__ double test___dsqrt_rn(double x) { // DEFAULT-LABEL: @test__fma_rn( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]] -// DEFAULT-NEXT: ret double [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// DEFAULT-NEXT: ret double [[TMP0]] // // FINITEONLY-LABEL: @test__fma_rn( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fma_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// FINITEONLY-NEXT: ret double [[TMP0]] // extern "C" __device__ double test__fma_rn(double x, double y, double z) { return __fma_rn(x, y, z); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
