[clang] [llvm] Reapply "DAGCombiner: Drop AllowFPOpFusion from visitFADDForFMACombine" (#221561) (PR #221567)
via cfe-commits
cfe-commits at lists.llvm.org
Sun Sep 6 06:03:31 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-arm
Author: Matt Arsenault (arsenm)
<details>
<summary>Changes</summary>
This reverts commit 502e51aa4df687807fbe51fa0b419baf65e9615f.
---
Patch is 457.36 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/221567.diff
24 Files Affected:
- (modified) clang/test/CodeGenCUDA/fp-contract.cu (+6-10)
- (modified) llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp (+3-3)
- (modified) llvm/test/CodeGen/AArch64/arm64-fma-combine-with-fpfusion.ll (+3-3)
- (modified) llvm/test/CodeGen/AArch64/fp-dp3.ll (+97-55)
- (modified) llvm/test/CodeGen/AArch64/neon-fma.ll (+13-13)
- (modified) llvm/test/CodeGen/AMDGPU/madak.ll (+388-471)
- (modified) llvm/test/CodeGen/ARM/fp16-instructions.ll (+10-10)
- (modified) llvm/test/CodeGen/Hexagon/float-amode.ll (+3-3)
- (modified) llvm/test/CodeGen/Hexagon/fmadd.ll (+3-3)
- (modified) llvm/test/CodeGen/Hexagon/fp_latency.ll (+7-7)
- (modified) llvm/test/CodeGen/LoongArch/double-fma.ll (+288-900)
- (modified) llvm/test/CodeGen/LoongArch/float-fma.ll (+288-900)
- (modified) llvm/test/CodeGen/LoongArch/lasx/fma-v4f64.ll (+372-1136)
- (modified) llvm/test/CodeGen/LoongArch/lasx/fma-v8f32.ll (+372-1136)
- (modified) llvm/test/CodeGen/LoongArch/lsx/fma-v2f64.ll (+178-564)
- (modified) llvm/test/CodeGen/LoongArch/lsx/fma-v4f32.ll (+178-564)
- (modified) llvm/test/CodeGen/PowerPC/a2-fp-basic.ll (+9-9)
- (modified) llvm/test/CodeGen/PowerPC/fma-assoc.ll (+122-125)
- (modified) llvm/test/CodeGen/PowerPC/fma-ext.ll (+16-16)
- (modified) llvm/test/CodeGen/PowerPC/fma.ll (+43-43)
- (modified) llvm/test/CodeGen/PowerPC/ppc440-fp-basic.ll (+9-9)
- (modified) llvm/test/CodeGen/PowerPC/vsx-fma-mutate-trivial-copy.ll (+4-4)
- (modified) llvm/test/CodeGen/SPARC/float-ua2007.ll (+18-18)
- (modified) llvm/test/CodeGen/Thumb2/mve-fmas.ll (+161-286)
``````````diff
diff --git a/clang/test/CodeGenCUDA/fp-contract.cu b/clang/test/CodeGenCUDA/fp-contract.cu
index 8e2942fc8104d..d98d5ca3f462b 100644
--- a/clang/test/CodeGenCUDA/fp-contract.cu
+++ b/clang/test/CodeGenCUDA/fp-contract.cu
@@ -252,21 +252,16 @@ __host__ __device__ float func2(float a, float b, float c) {
// AMD-OPT-OFF-NEXT: v_add_f32_e32
// AMD-OPT-OFF-NEXT: s_setpc_b64
-// Test multiply/add in the different statements, which is forced
-// to be compiled with fp contract on. fmul/fadd without contract
-// flags are emitted in IR. In nvptx, they are emitted as FMA in
-// fp-contract is fast but not on, as nvptx backend uses the same
-// fp fuse option as front end, whereas fast fp fuse option in
-// backend fuses fadd/fmul disregarding contract flag. In amdgcn
-// they are not fused as amdgcn always use standard fp fusion
-// option which respects contract flag.
+// Test multiply/add in the different statements, which is forced to be compiled
+// with fp contract on. fmul/fadd without contract flags are emitted in IR.
__host__ __device__ float func3(float a, float b, float c) {
#pragma clang fp contract(on)
float t = b * c;
return t + a;
}
// COMMON-LABEL: _Z5func3fff
-// NV-OPT-FAST: fma.rn.f32
+// NV-OPT-FAST: mul.rn.f32
+// NV-OPT-FAST: add.rn.f32
// NV-OPT-FAST-NEXT: st.param.b32
// NV-OPT-FASTSTD: mul.rn.f32
// NV-OPT-FASTSTD: add.rn.f32
@@ -285,7 +280,8 @@ __host__ __device__ float func2(float a, float b, float c) {
// AMD-OPT-OFF-IR: fmul float
// AMD-OPT-OFF-IR: fadd float
-// AMD-OPT-FAST: v_fmac_f32_e32
+// AMD-OPT-FAST: v_mul_f32_e32
+// AMD-OPT-FAST-NEXT: v_add_f32_e32
// AMD-OPT-FAST-NEXT: s_setpc_b64
// AMD-OPT-FASTSTD: v_mul_f32_e32
// AMD-OPT-FASTSTD-NEXT: v_add_f32_e32
diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
index 70e23360ed9ea..82d4921ae405b 100644
--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -18679,7 +18679,6 @@ SDValue DAGCombiner::visitFADDForFMACombine(SDNode *N) {
SDValue N1 = N->getOperand(1);
EVT VT = N->getValueType(0);
SDLoc SL(N);
- const TargetOptions &Options = DAG.getTarget().Options;
// Floating-point multiply-add with intermediate rounding.
bool HasFMAD = (LegalOperations && TLI.isFMADLegal(DAG, N));
@@ -18693,8 +18692,9 @@ SDValue DAGCombiner::visitFADDForFMACombine(SDNode *N) {
if (!HasFMAD && !HasFMA)
return SDValue();
- bool AllowFusionGlobally =
- Options.AllowFPOpFusion == FPOpFusion::Fast || HasFMAD;
+ // FMAD (with intermediate rounding) is always safe to form; FMA requires the
+ // contract fast-math flag.
+ bool AllowFusionGlobally = HasFMAD;
// If the addition is not contractable, do not combine.
if (!AllowFusionGlobally && !N->getFlags().hasAllowContract())
return SDValue();
diff --git a/llvm/test/CodeGen/AArch64/arm64-fma-combine-with-fpfusion.ll b/llvm/test/CodeGen/AArch64/arm64-fma-combine-with-fpfusion.ll
index 0ddcdccca4ec4..daa9d38db78e3 100644
--- a/llvm/test/CodeGen/AArch64/arm64-fma-combine-with-fpfusion.ll
+++ b/llvm/test/CodeGen/AArch64/arm64-fma-combine-with-fpfusion.ll
@@ -1,10 +1,10 @@
-; RUN: llc < %s -mtriple=aarch64-linux-gnu -fp-contract=fast | FileCheck %s
+; RUN: llc < %s -mtriple=aarch64-linux-gnu | FileCheck %s
define float @mul_add(float %a, float %b, float %c) local_unnamed_addr #0 {
; CHECK-LABEL: %entry
; CHECK: fmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
entry:
- %mul = fmul float %a, %b
- %add = fadd float %mul, %c
+ %mul = fmul contract float %a, %b
+ %add = fadd contract float %mul, %c
ret float %add
}
diff --git a/llvm/test/CodeGen/AArch64/fp-dp3.ll b/llvm/test/CodeGen/AArch64/fp-dp3.ll
index 10f88fdbbe962..e5736eeb5c275 100644
--- a/llvm/test/CodeGen/AArch64/fp-dp3.ll
+++ b/llvm/test/CodeGen/AArch64/fp-dp3.ll
@@ -1,137 +1,177 @@
-; RUN: llc -verify-machineinstrs -o - %s -mtriple=aarch64-none-linux-gnu -fp-contract=fast | FileCheck %s
-; RUN: llc -verify-machineinstrs -o - %s -mtriple=arm64-apple-ios7.0 | FileCheck %s -check-prefix=CHECK-NOFAST
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=aarch64-none-linux < %s | FileCheck %s
declare float @llvm.fma.f32(float, float, float)
declare double @llvm.fma.f64(double, double, double)
define float @test_fmadd(float %a, float %b, float %c) {
; CHECK-LABEL: test_fmadd:
-; CHECK-NOFAST-LABEL: test_fmadd:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmadd s0, s0, s1, s2
+; CHECK-NEXT: ret
%val = call float @llvm.fma.f32(float %a, float %b, float %c)
-; CHECK: fmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %val
}
define float @test_fmsub(float %a, float %b, float %c) {
; CHECK-LABEL: test_fmsub:
-; CHECK-NOFAST-LABEL: test_fmsub:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmsub s0, s0, s1, s2
+; CHECK-NEXT: ret
%nega = fsub float -0.0, %a
%val = call float @llvm.fma.f32(float %nega, float %b, float %c)
-; CHECK: fmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %val
}
define float @test_fnmadd(float %a, float %b, float %c) {
; CHECK-LABEL: test_fnmadd:
-; CHECK-NOFAST-LABEL: test_fnmadd:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fnmadd s0, s0, s1, s2
+; CHECK-NEXT: ret
%nega = fsub float -0.0, %a
%negc = fsub float -0.0, %c
%val = call float @llvm.fma.f32(float %nega, float %b, float %negc)
-; CHECK: fnmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fnmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %val
}
define float @test_fnmsub(float %a, float %b, float %c) {
; CHECK-LABEL: test_fnmsub:
-; CHECK-NOFAST-LABEL: test_fnmsub:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fnmsub s0, s0, s1, s2
+; CHECK-NEXT: ret
%negc = fsub float -0.0, %c
%val = call float @llvm.fma.f32(float %a, float %b, float %negc)
-; CHECK: fnmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fnmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %val
}
define double @testd_fmadd(double %a, double %b, double %c) {
; CHECK-LABEL: testd_fmadd:
-; CHECK-NOFAST-LABEL: testd_fmadd:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmadd d0, d0, d1, d2
+; CHECK-NEXT: ret
%val = call double @llvm.fma.f64(double %a, double %b, double %c)
-; CHECK: fmadd {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
-; CHECK-NOFAST: fmadd {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
ret double %val
}
define double @testd_fmsub(double %a, double %b, double %c) {
; CHECK-LABEL: testd_fmsub:
-; CHECK-NOFAST-LABEL: testd_fmsub:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmsub d0, d0, d1, d2
+; CHECK-NEXT: ret
%nega = fsub double -0.0, %a
%val = call double @llvm.fma.f64(double %nega, double %b, double %c)
-; CHECK: fmsub {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
-; CHECK-NOFAST: fmsub {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
ret double %val
}
define double @testd_fnmadd(double %a, double %b, double %c) {
; CHECK-LABEL: testd_fnmadd:
-; CHECK-NOFAST-LABEL: testd_fnmadd:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fnmadd d0, d0, d1, d2
+; CHECK-NEXT: ret
%nega = fsub double -0.0, %a
%negc = fsub double -0.0, %c
%val = call double @llvm.fma.f64(double %nega, double %b, double %negc)
-; CHECK: fnmadd {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
-; CHECK-NOFAST: fnmadd {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
ret double %val
}
define double @testd_fnmsub(double %a, double %b, double %c) {
; CHECK-LABEL: testd_fnmsub:
-; CHECK-NOFAST-LABEL: testd_fnmsub:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fnmsub d0, d0, d1, d2
+; CHECK-NEXT: ret
%negc = fsub double -0.0, %c
%val = call double @llvm.fma.f64(double %a, double %b, double %negc)
-; CHECK: fnmsub {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
-; CHECK-NOFAST: fnmsub {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}, {{d[0-9]+}}
ret double %val
}
define float @test_fmadd_unfused(float %a, float %b, float %c) {
; CHECK-LABEL: test_fmadd_unfused:
-; CHECK-NOFAST-LABEL: test_fmadd_unfused:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmul s1, s1, s2
+; CHECK-NEXT: fadd s0, s0, s1
+; CHECK-NEXT: ret
%prod = fmul float %b, %c
%sum = fadd float %a, %prod
-; CHECK: fmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST-NOT: fmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fmul {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %sum
}
define float @test_fmsub_unfused(float %a, float %b, float %c) {
; CHECK-LABEL: test_fmsub_unfused:
-; CHECK-NOFAST-LABEL: test_fmsub_unfused:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmul s1, s1, s2
+; CHECK-NEXT: fsub s0, s0, s1
+; CHECK-NEXT: ret
%prod = fmul float %b, %c
%diff = fsub float %a, %prod
-; CHECK: fmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST-NOT: fmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fmul {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %diff
}
define float @test_fnmadd_unfused(float %a, float %b, float %c) {
; CHECK-LABEL: test_fnmadd_unfused:
-; CHECK-NOFAST-LABEL: test_fnmadd_unfused:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fneg s0, s0
+; CHECK-NEXT: fmul s1, s1, s2
+; CHECK-NEXT: fsub s0, s0, s1
+; CHECK-NEXT: ret
%nega = fsub float -0.0, %a
%prod = fmul float %b, %c
%diff = fsub float %nega, %prod
-; CHECK: fnmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST-NOT: fnmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fmul {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: ret
ret float %diff
}
define float @test_fnmsub_unfused(float %a, float %b, float %c) {
; CHECK-LABEL: test_fnmsub_unfused:
-; CHECK-NOFAST-LABEL: test_fnmsub_unfused:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmul s1, s1, s2
+; CHECK-NEXT: fsub s0, s1, s0
+; CHECK-NEXT: ret
%nega = fsub float -0.0, %a
%prod = fmul float %b, %c
%sum = fadd float %nega, %prod
-; CHECK: fnmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST-NOT: fnmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fmul {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK-NOFAST: fsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
+ ret float %sum
+}
+
+; With the contract flag the multiply and add/subtract fuse to a single op.
+
+define float @test_fmadd_contract(float %a, float %b, float %c) {
+; CHECK-LABEL: test_fmadd_contract:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmadd s0, s1, s2, s0
+; CHECK-NEXT: ret
+ %prod = fmul contract float %b, %c
+ %sum = fadd contract float %a, %prod
+ ret float %sum
+}
+
+define float @test_fmsub_contract(float %a, float %b, float %c) {
+; CHECK-LABEL: test_fmsub_contract:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmsub s0, s1, s2, s0
+; CHECK-NEXT: ret
+ %prod = fmul contract float %b, %c
+ %diff = fsub contract float %a, %prod
+ ret float %diff
+}
+
+define float @test_fnmadd_contract(float %a, float %b, float %c) {
+; CHECK-LABEL: test_fnmadd_contract:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fnmadd s0, s1, s2, s0
+; CHECK-NEXT: ret
+ %nega = fsub float -0.0, %a
+ %prod = fmul contract float %b, %c
+ %diff = fsub contract float %nega, %prod
+ ret float %diff
+}
+
+define float @test_fnmsub_contract(float %a, float %b, float %c) {
+; CHECK-LABEL: test_fnmsub_contract:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fnmsub s0, s1, s2, s0
+; CHECK-NEXT: ret
+ %nega = fsub float -0.0, %a
+ %prod = fmul contract float %b, %c
+ %sum = fadd contract float %nega, %prod
ret float %sum
}
@@ -139,25 +179,27 @@ define float @test_fnmsub_unfused(float %a, float %b, float %c) {
define float @test_fmadd_unfused_su(float %a, float %b, float %c) {
; CHECK-LABEL: test_fmadd_unfused_su:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmul s1, s1, s2
+; CHECK-NEXT: fadd s0, s0, s1
+; CHECK-NEXT: fadd s0, s0, s1
+; CHECK-NEXT: ret
%prod = fmul float %b, %c
%sum = fadd float %a, %prod
%res = fadd float %sum, %prod
-; CHECK-NOT: fmadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK: fmul {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK: fadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK: fadd {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %res
}
define float @test_fmsub_unfused_su(float %a, float %b, float %c) {
; CHECK-LABEL: test_fmsub_unfused_su:
+; CHECK: // %bb.0:
+; CHECK-NEXT: fmul s1, s1, s2
+; CHECK-NEXT: fsub s0, s0, s1
+; CHECK-NEXT: fsub s0, s0, s1
+; CHECK-NEXT: ret
%prod = fmul float %b, %c
%diff = fsub float %a, %prod
%res = fsub float %diff, %prod
-; CHECK-NOT: fmsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK: fmul {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK: fsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
-; CHECK: fsub {{s[0-9]+}}, {{s[0-9]+}}, {{s[0-9]+}}
ret float %res
}
diff --git a/llvm/test/CodeGen/AArch64/neon-fma.ll b/llvm/test/CodeGen/AArch64/neon-fma.ll
index af70302ca939c..25aeb6be65939 100644
--- a/llvm/test/CodeGen/AArch64/neon-fma.ll
+++ b/llvm/test/CodeGen/AArch64/neon-fma.ll
@@ -1,45 +1,45 @@
-; RUN: llc < %s -verify-machineinstrs -mtriple=aarch64-none-linux-gnu -mattr=+neon -fp-contract=fast | FileCheck %s
+; RUN: llc < %s -verify-machineinstrs -mtriple=aarch64-none-linux-gnu -mattr=+neon | FileCheck %s
define <2 x float> @fmla2xfloat(<2 x float> %A, <2 x float> %B, <2 x float> %C) {
;CHECK: fmla {{v[0-9]+}}.2s, {{v[0-9]+}}.2s, {{v[0-9]+}}.2s
- %tmp1 = fmul <2 x float> %A, %B;
- %tmp2 = fadd <2 x float> %C, %tmp1;
+ %tmp1 = fmul contract <2 x float> %A, %B;
+ %tmp2 = fadd contract <2 x float> %C, %tmp1;
ret <2 x float> %tmp2
}
define <4 x float> @fmla4xfloat(<4 x float> %A, <4 x float> %B, <4 x float> %C) {
;CHECK: fmla {{v[0-9]+}}.4s, {{v[0-9]+}}.4s, {{v[0-9]+}}.4s
- %tmp1 = fmul <4 x float> %A, %B;
- %tmp2 = fadd <4 x float> %C, %tmp1;
+ %tmp1 = fmul contract <4 x float> %A, %B;
+ %tmp2 = fadd contract <4 x float> %C, %tmp1;
ret <4 x float> %tmp2
}
define <2 x double> @fmla2xdouble(<2 x double> %A, <2 x double> %B, <2 x double> %C) {
;CHECK: fmla {{v[0-9]+}}.2d, {{v[0-9]+}}.2d, {{v[0-9]+}}.2d
- %tmp1 = fmul <2 x double> %A, %B;
- %tmp2 = fadd <2 x double> %C, %tmp1;
+ %tmp1 = fmul contract <2 x double> %A, %B;
+ %tmp2 = fadd contract <2 x double> %C, %tmp1;
ret <2 x double> %tmp2
}
define <2 x float> @fmls2xfloat(<2 x float> %A, <2 x float> %B, <2 x float> %C) {
;CHECK: fmls {{v[0-9]+}}.2s, {{v[0-9]+}}.2s, {{v[0-9]+}}.2s
- %tmp1 = fmul <2 x float> %A, %B;
- %tmp2 = fsub <2 x float> %C, %tmp1;
+ %tmp1 = fmul contract <2 x float> %A, %B;
+ %tmp2 = fsub contract <2 x float> %C, %tmp1;
ret <2 x float> %tmp2
}
define <4 x float> @fmls4xfloat(<4 x float> %A, <4 x float> %B, <4 x float> %C) {
;CHECK: fmls {{v[0-9]+}}.4s, {{v[0-9]+}}.4s, {{v[0-9]+}}.4s
- %tmp1 = fmul <4 x float> %A, %B;
- %tmp2 = fsub <4 x float> %C, %tmp1;
+ %tmp1 = fmul contract <4 x float> %A, %B;
+ %tmp2 = fsub contract <4 x float> %C, %tmp1;
ret <4 x float> %tmp2
}
define <2 x double> @fmls2xdouble(<2 x double> %A, <2 x double> %B, <2 x double> %C) {
;CHECK: fmls {{v[0-9]+}}.2d, {{v[0-9]+}}.2d, {{v[0-9]+}}.2d
- %tmp1 = fmul <2 x double> %A, %B;
- %tmp2 = fsub <2 x double> %C, %tmp1;
+ %tmp1 = fmul contract <2 x double> %A, %B;
+ %tmp2 = fsub contract <2 x double> %C, %tmp1;
ret <2 x double> %tmp2
}
diff --git a/llvm/test/CodeGen/AMDGPU/madak.ll b/llvm/test/CodeGen/AMDGPU/madak.ll
index 211b93e9ace29..4e640573c0193 100644
--- a/llvm/test/CodeGen/AMDGPU/madak.ll
+++ b/llvm/test/CodeGen/AMDGPU/madak.ll
@@ -2,13 +2,10 @@
; RUN: llc -mtriple=amdgpu6.00 < %s | FileCheck -check-prefixes=GFX6 %s
; RUN: llc -mtriple=amdgpu8.02 < %s | FileCheck -check-prefixes=GFX8 %s
; RUN: llc -mtriple=amdgpu9.00 < %s | FileCheck -check-prefixes=GFX9 %s
+; RUN: llc -mtriple=amdgpu9.42 < %s | FileCheck -check-prefixes=GFX942 %s
; RUN: llc -mtriple=amdgpu10.10 < %s | FileCheck -check-prefixes=GFX10-MAD %s
; RUN: llc -mtriple=amdgpu11.00 < %s | FileCheck -check-prefixes=GFX11-MAD %s
-; RUN: llc -mtriple=amdgpu9.42 -fp-contract=fast < %s | FileCheck -check-prefixes=GFX942-FMA %s
-; RUN: llc -mtriple=amdgpu10.10 -fp-contract=fast < %s | FileCheck -check-prefixes=GFX10-FMA %s
-; RUN: llc -mtriple=amdgpu11.00 -fp-contract=fast < %s | FileCheck -check-prefixes=GFX11-FMA %s
-
declare i32 @llvm.amdgcn.workitem.id.x() nounwind readnone
declare float @llvm.fabs.f32(float) nounwind readnone
@@ -67,6 +64,21 @@ define amdgpu_kernel void @madak_f32(ptr addrspace(1) noalias %out, ptr addrspac
; GFX9-NEXT: global_store_dword v0, v1, s[0:1]
; GFX9-NEXT: s_endpgm
;
+; GFX942-LABEL: madak_f32:
+; GFX942: ; %bb.0:
+; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
+; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
+; GFX942-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; GFX942-NEXT: v_lshlrev_b32_e32 v0, 2, v0
+; GFX942-NEXT: s_waitcnt lgkmcnt(0)
+; GFX942-NEXT: global_load_dword v1, v0, s[2:3]
+; GFX942-NEXT: global_load_dword v2, v0, s[6:7]
+; GFX942-NEXT: s_waitcnt vmcnt(0)
+; GFX942-NEXT: v_mul_f32_e32 v1, v1, v2
+; GFX942-NEXT: v_add_f32_e32 v1, 0x41200000, v1
+; GFX942-NEXT: global_store_dword v0, v1, s[0:1]
+; GFX942-NEXT: s_endpgm
+;
; GFX10-MAD-LABEL: madak_f32:
; GFX10-MAD: ; %bb.0:
; GFX10-MAD-NEXT: s_clause 0x1
@@ -100,52 +112,6 @@ define amdgpu_kernel void @madak_f32(ptr addrspace(1) noalias %out, ptr addrspac
; GFX11-MAD-NEXT: v_add_f32_e32 v1, 0x41200000, v1
; GFX11-MAD-NEXT: global_store_b32 v0, v1, s[0:1]
; GFX11-MAD-NEXT: s_endpgm
-;
-; GFX942-FMA-LABEL: madak_f32:
-; GFX942-FMA: ; %bb.0:
-; GFX942-FMA-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
-; GFX942-FMA-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
-; GFX942-FMA-NEXT: v_and_b32_e32 v0, 0x3ff, v0
-; GFX942-FMA-NEXT: v_lshlrev_b32_e32 v0, 2, v0
-; GFX942-FMA-NEXT: s_waitcnt lgkmcnt(0)
-; GFX942-FMA-NEXT: global_load_dword v1, v0, s[2:3]
-; GFX942-FMA-NEXT: global_load_dword v2, v0, s[6:7]
-; GFX942-FMA-NEXT: s_waitcnt vmcnt(0)
-; GFX942-FMA-NEXT: v_fmaak_f32 v1, v1, v2, 0x41200000
-; GFX942-FMA-NEXT: global_store_dword v0, v1, s[0:1]
-; GFX942-FMA-NEXT: s_endpgm
-;
-; GFX10-FMA-LABEL: madak_f32:
-; GFX10-FMA: ; %bb.0:
-; GFX10-FMA-NEXT: s_clause 0x1
-; GFX10-FMA-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
-; GFX10-FMA-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
-; GFX10-FMA-NEXT: v_lshlrev_b32_e32 v0, 2, v0
-; GFX10-FMA-NEXT: s_waitcnt lgkmcnt(0)
-; GFX10-FMA-NEXT: s_clause 0x1
-; GFX10-FMA-NEXT: global_load_dword v1, v0, s[2:3]
-; GFX10-FMA-NEXT: global_load_dword v2, v0, s[6:7]
-; GFX10-FMA-NEXT: s_waitcnt vmcnt(0)
-; GFX10-FMA-NEXT: v_fmaak_f32 v1, v1, v2, 0x41200000
-; GFX10-FMA-NEXT: global_store_dword v0, v1, s[0:1]
-; GFX10-FMA-NEXT: s_endpgm
-;
-; GFX11-FMA-LABEL: madak_f32:
-; GFX11-FMA: ; %bb.0:
-; GFX11-FMA-NEXT: s_clause 0x1
-; GFX11-FMA-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
-; GFX11-FMA-NEXT: s_load_b64 s[4:5], s[4:5], 0x34
-; GFX11-FMA-NEXT: v_and_b32_e32 v0, 0x3ff, v0
-; GFX11-FMA-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; GFX11-FMA-NEXT: v_lshlrev_b32_e32 v0, 2, v0
-; GFX11-FMA-NEXT: s_waitcnt lgkmcnt(0)
-; GFX11-FMA-NEXT: s_clause 0x1
-; GFX11-FMA-NEXT: global_load_b32 v1, v0, s[2:3]
-; GFX11-FMA-NEXT: global_load_b32 v2, v0, s[4:5]
-; GFX11-FMA-NEXT: s_waitcnt vmcnt(0)
-; GFX11-FMA-NEXT: v_fmaak_f32 v1, v1, v2, 0x41200000
-; GFX11-FMA-NEXT: global_store_b32 v0, v1, s[0:1]
-; GFX11-FMA-NEXT: s_endpgm
%tid = tail call i32 @llvm.amdgcn.workitem.id.x() nounwind readnone
%in.a.gep = getelementptr float, ptr addrspace(1) %in.a, i32 %tid
%in.b.gep = getelementptr float, ptr addrspace(1) %in.b, i32 %tid
@@ -239,6 +205,28 @@ define amdgpu_kernel void @madak_2_use_f32(ptr addrspace(1) noalias %out, ptr ad
; GFX9-NEXT: s_waitc...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/221567
More information about the cfe-commits
mailing list