[llvm] [AMDGPU] Add RewriteMFMAForm scheduling stage tests (PR #196008)

[Matt Arsenault via llvm-commits]

================
@@ -0,0 +1,332 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 \
+; RUN:     -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN:     < %s | FileCheck %s
+;
+; Test: RewriteMFMAFormStage — v7_slice pattern (gfx950, mfma.f32.16x16x32.f16)
+;
+; Distilled from v7_slice.llir (Triton matmul kernel, gfx950).
+;
+; Key structural features preserved from v7_slice:
+;   (1) Loop body (%loop) accumulates via loop-carried SCALAR float phis,
+;       NOT <4 x float> phis. Each MFMA acc is built by insertelement from
+;       4 individual float phis then used as src2/dst in MFMAs within the loop.
+;   (2) Epilogue (%epilogue) ALSO contains MFMAs: the final K-tile iteration
+;       is peeled. Epilogue MFMAs take scalar float phis from %loop as acc,
+;       then produce results consumed by fptrunc+store — Case 2.
+;   (3) Scalar float phis initialized to 0.0 in entry (non-MAI def) — Case 3.
+;
+; CFG: entry -> loop (back-edge) -> epilogue -> ret
+;
+; Pressure design (gfx950, 256 ArchVGPR limit):
+;   16 loop-carried scalar float phis (4 chains x 4 elems) = 16 VGPRs.
+;   4 MFMA chains of 2, each acc = 4 float phis = 16 VGPRs in loop.
+;   4 x <8 x half> operands (a0, a1, b0, b1) = 16 VGPRs.
+;   8 x <32 x float> loop-carried vec carriers (vc0-vc7) = 256 VGPRs.
+;   Peak ArchVGPR = 256 + 16 + 16 + 16 = 304 > 256 -> RegionsWithExcessArchVGPR.
+;   After rewrite: scalar phis=16 VGPR, <8 x half>=16 VGPR, vec=256 VGPR, MFMA acc->AGPR.
+;   getRewriteCost() < 0 -> rewrite() fires.
+;
+; Expected:
+;   Case 3: v_accvgpr_write inserted after insertelement defs of MFMA acc
+;           (entry zeroinitializer -> scalar float phi -> insertelement = non-MAI)
+;   Case 2: v_accvgpr_read before fptrunc in epilogue block
+;   MFMA opcodes: v_mfma_f32_16x16x32_f16 vgprcd -> agprcd form
+
----------------
arsenm wrote:

Same comment, this looks like overly verbose AI description 

https://github.com/llvm/llvm-project/pull/196008