AMDGPU: Add baseline test for vgpr mfma with copied-from AGPR #153020

arsenm · 2025-08-11T14:40:02Z

We currently handle the case where an MFMA is copied to an AGPR,
but not the case where the MFMA is copying from an AGPR.

arsenm · 2025-08-11T14:40:23Z

This stack of pull requests is managed by Graphite. Learn more about stacking.

llvmbot · 2025-08-11T14:40:38Z

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

We currently handle the case where an MFMA is copied to an AGPR,
but not the case where the MFMA is copying from an AGPR.

Patch is 34.43 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/153020.diff

2 Files Affected:

(added) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir (+261)
(modified) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll (+387)

diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir
new file mode 100644
index 0000000000000..7fdc8c0d8019b
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir
@@ -0,0 +1,261 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -run-pass=greedy,amdgpu-rewrite-agpr-copy-mfma -verify-machineinstrs -o - %s | FileCheck %s
+
+---
+name:  test_rewrite_mfma_copy_from_agpr_physreg
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $agpr0_agpr1
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_physreg
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $agpr0_agpr1
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_64_align2 = COPY $agpr0_agpr1
+    ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:vreg_64_align2 = COPY $agpr0_agpr1
+    %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0, implicit $mode, implicit $exec
+    GLOBAL_STORE_DWORDX2 %0, %4, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    SI_RETURN
+...
+
+---
+name:  test_rewrite_mfma_copy_from_agpr_unrewritable_use
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_unrewritable_use
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+    ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3866633 /* reguse:VReg_64_Align2 */, [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]]
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    %4:vreg_128_align2 = COPY %3
+    %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+    INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3866633 /* reguse:VReg_64_Align2 */, %5
+    SI_RETURN
+...
+
+---
+name:  test_rewrite_mfma_copy_from_agpr_src2_subreg_use
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_src2_subreg_use
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+    ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    %4:vreg_128_align2 = COPY %3
+    %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+    GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    SI_RETURN
+...
+
+---
+name:  test_rewrite_mfma_copy_from_agpr_vdst_subreg_use
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit $exec :: (store (s128), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    %4:vreg_128_align2 = COPY %3
+    %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+    GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1)
+    SI_RETURN
+...
+
+# A-to-V copy is performed subregister at a time instead.
+---
+name:  test_rewrite_mfma_copy_from_agpr_split_copy
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_split_copy
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+    ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]].sub0
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub1:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]].sub1
+    ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1)
+    undef %4.sub0:vreg_64_align2 = COPY %3.sub0
+    %4.sub1:vreg_64_align2 = COPY %3.sub1
+    %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4, 0, 0, 0, implicit $mode, implicit $exec
+    GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    SI_RETURN
+...
+
+---
+name:  test_rewrite_mfma_copy_from_agpr_copyback
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_copyback
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]]
+    ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_e64_:%[0-9]+]]:areg_64_align2 = V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: [[COPY4:%[0-9]+]]:areg_64_align2 = COPY [[V_MFMA_F64_4X4X4F64_e64_]]
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY4]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1)
+    %4:vreg_64_align2 = COPY %3
+    %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4, 0, 0, 0, implicit $mode, implicit $exec
+    %6:areg_64_align2 = COPY %5
+    GLOBAL_STORE_DWORDX2 %0, %6, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    SI_RETURN
+...
+
+# There is a read of the coy-from-agpr in the dst operand of the MFMA.
+---
+name:  test_rewrite_mfma_copy_from_agpr_vdst_subreg_use_imm_src2
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use_imm_src2
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+    ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit $exec :: (store (s128), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    %4:vreg_128_align2 = COPY %3
+    %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec
+    GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1)
+    SI_RETURN
+...
+
+# Degenerate case. Copy from AGPR to VGPR is dead undef subreg def
+---
+name:  test_rewrite_mfma_copy_from_agpr_undef_vdst_subreg_use_imm_src2
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_undef_vdst_subreg_use_imm_src2
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    ; CHECK-NEXT: dead [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+    ; CHECK-NEXT: undef [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s128), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+    %4:vreg_128_align2 = COPY %3
+    undef %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec
+    GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1)
+    SI_RETURN
+...
+
+# Degenerate case. Copy from AGPR to VGPR is dead, but same register
+# is redefined as whole register.
+---
+name:  test_rewrite_mfma_copy_from_agpr_to_vdst_def_imm_src2
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+    ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_to_vdst_def_imm_src2
+    ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+    ; CHECK-NEXT: {{  $}}
+    ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+    ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+    ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+    ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+    ; CHECK-NEXT: dead [[COPY3:%[0-9]+]]:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]]
+    ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec
+    ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    ; CHECK-NEXT: SI_RETURN
+    %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+    %1:av_64_align2 = COPY $vgpr0_vgpr1
+    %2:av_64_align2 = COPY $vgpr2_vgpr3
+    %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1)
+    %4:vreg_64_align2 = COPY %3
+    %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec
+    GLOBAL_STORE_DWORDX2 %0, %4, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+    SI_RETURN
+...
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
index 58ec41e3928bd..81613f69c982b 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
@@ -522,7 +522,394 @@ define void @test_rewrite_mfma_subreg_insert2(double %arg0, double %arg1, ptr ad
   ret void
 }
 
+define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class(ptr addrspace(1) %arg0, ptr addrspace(1) %arg1) #0 {
+; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class:
+; CHECK:       ; %bb.0:
+; CHECK-NEXT:    ;;#ASMSTART
+; CHECK-NEXT:    ; def a[0:31]
+; CHECK-NEXT:    ;;#ASMEND
+; CHECK-NEXT:    v_accvgpr_write_b32 a32, v0
+; CHECK-NEXT:    v_accvgpr_read_b32 v63, a31
+; CHECK-NEXT:    v_accvgpr_read_b32 v62, a30
+; CHECK-NEXT:    v_accvgpr_read_b32 v61, a29
+; CHECK-NEXT:    v_accvgpr_read_b32 v60, a28
+; CHECK-NEXT:    v_accvgpr_read_b32 v59, a27
+; CHECK-NEXT:    v_accvgpr_read_b32 v58, a26
+; CHECK-NEXT:    v_accvgpr_read_b32 v57, a25
+; CHECK-NEXT:    v_accvgpr_read_b32 v56, a24
+; CHECK-NEXT:    v_accvgpr_read_b32 v55, a23
+; CHECK-NEXT:    v_accvgpr_read_b32 v54, a22
+; CHECK-NEXT:    v_accvgpr_read_b32 v53, a21
+; CHECK-NEXT:    v_accvgpr_read_b32 v52, a20
+; CHECK-NEXT:    v_accvgpr_read_b32 v51, a19
+; CHECK-NEXT:    v_accvgpr_read_b32 v50, a18
+; CHECK-NEXT:    v_accvgpr_read_b32 v49, a17
+; CHECK-NEXT:    v_accvgpr_read_b32 v48, a16
+; CHECK-NEXT:    v_accvgpr_read_b32 v47, a15
+; CHECK-NEXT:    v_accvgpr_read_b32 v46, a14
+; CHECK-NEXT:    v_accvgpr_read_b32 v45, a13
+; CHECK-NEXT:    v_accvgpr_read_b32 v44, a12
+; CHECK-NEXT:    v_accvgpr_read_b32 v43, a11
+; CHECK-NEXT:    v_accvgpr_read_b32 v42, a10
+; CHECK-NEXT:    v_accvgpr_read_b32 v41, a9
+; CHECK-NEXT:    v_accvgpr_read_b32 v40, a8
+; CHECK-NEXT:    v_accvgpr_read_b32 v39, a7
+; CHECK-NEXT:    v_accvgpr_read_b32 v38, a6
+; CHECK-NEXT:    v_accvgpr_read_b32 v37, a5
+; CHECK-NEXT:    v_accvgpr_read_b32 v36, a4
+; CHECK-NEXT:    v_accvgpr_read_b32 v35, a3
+; CHECK-NEXT:    v_accvgpr_read_b32 v34, a2
+; CHECK-NEXT:    v_accvgpr_read_b32 v33, a1
+; CHECK-NEXT:    v_accvgpr_read_b32 v32, a0
+; CHECK-NEXT:    v_accvgpr_write_b32 a0, 2.0
+; CHECK-NEXT:    v_accvgpr_write_b32 a1, 4.0
+; CHECK-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x0
+; CHECK-NEXT:    s_nop 0
+; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], a0, a1, v[32:63]
+; CHECK-NEXT:    v_accvgpr_write_b32 a0, v32
+; CHECK-NEXT:    v_accvgpr_write_b32 a1, v33
+; CHECK-NEXT:    v_accvgpr_write_b32 a2, v34
+; CHECK-NEXT:    v_accvgpr_write_b32 a3, v35
+; CHECK-NEXT:    v_accvgpr_write_b32 a4, v36
+; CHECK-NEXT:    v_accvgpr_write_b32 a5, v37
+; CHECK-NEXT:    v_accvgpr_write_b32 a6, v38
+; CHECK-NEXT:    v_accvgpr_write_b32 a7, v39
+; CHECK-NEXT:    v_accvgpr_write_b32 a8, v40
+; CHECK-NEXT:    v_accvgpr_write_b32 a9, v41
+; CHECK-NEXT:    v_accvgpr_write_b32 a10, v42
+; CHECK-NEXT:    v_accvgpr_write_b32 a11, v43
+; CHECK-NEXT:    v_accvgpr_write_b32 a12, v44
+; CHECK-NEXT:    v_accvgpr_write_b32 a13, v45
+; CHECK-NEXT:    v_accvgpr_write_b32 a14, v46
+; CHECK-NEXT:    v_accvgpr_write_b32 a15, v47
+; CHECK-NEXT:    v_accvgpr_write_b32 a16, v48
+; CHECK-NEXT:    v_accvgpr_write_b32 a17, v49
+; CHECK-NEXT:    v_accvgpr_write_b32 a18, v50
+; CHECK-NEXT:    v_accvgpr_write_b32 a19, v51
+; CHECK-NEXT:    v_accvgpr_write_b32 a20, v52
+; CHECK-NEXT:    v_accvgpr_write_b32 a21, v53
+; CHECK-NEXT:    v_accvgpr_write_b32 a22, v54
+; CHECK-NEXT:    v_accvgpr_write_b32 a23, v55
+; CHECK-NEXT:    v_accvgpr_write_b32 a24, v56
+; CHECK-NEXT:    v_accvgpr_write_b32 a25, v57
+; CHECK-NEXT:    v_accvgpr_write_b32 a26, v58
+; CHECK-NEXT:    v_accvgpr_write_b32 a27, v59
+; CHECK-NEXT:    v_accvgpr_write_b32 a28, v60
+; CHECK-NEXT:    v_accvgpr_write_b32 a29, v61
+; CHECK-NEXT:    v_accvgpr_write_b32 a30, v62
+; CHECK-NEXT:    v_accvgpr_write_b32 a31, v63
+; CHECK-NEXT:    v_accvgpr_read_b32 v32, a32
+; CHECK-NEXT:    v_mov_b32_e32 v33, 0x41000000
+; CHECK-NEXT:    v_and_b32_e32 v32, 0x3ff, v32
+; CHECK-NEXT:    v_lshlrev_b32_e32 v32, 7, v32
+; CHECK-NEXT:    s_waitcnt lgkmcnt(0)
+; CHECK-NEXT:    global_store_dwordx4 v32, v[28:31], s[0:1] offset:112
+; CHECK-NEXT:    global_store_dwordx4 v32, v[24:27], s[0:1] offset:96
+; CHECK-NEXT:    global_store_dwordx4 v32, v[20:23], s[0:1] offset:80
+; CHECK-NEXT:    global_store_dwordx4 v32, v[16:19], s[0:1] offset:64
+; CHECK-NEXT:    global_store_dwordx4 v32, v[12:15], s[0:1] offset:48
+; CHECK-NEXT:    global_store_dwordx4 v32, v[8:11], s[0:1] offset:32
+; CHECK-NEXT:    global_store_dwordx4 v32, v[4:7], s[0:1] offset:16
+; CHECK-NEXT:    global_store_dwordx4 v32, v[0:3], s[0:1]
+; CHECK-NEXT:    v_mov_b32_e32 v34, 0x41800000
+; CHECK-NEXT:    s_nop 0
+; CHECK-NEXT:    v_accvgpr_read_b32 v0, a0
+; CHECK-NEXT:    v_accvgpr_read_b32 v1, a1
+; CHECK-NEXT:    v_accvgpr_read_b3...
[truncated]

llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir

This should address the case where the result isn't fully used, resulting in partial copy bundles from the MFMA result.

We currently handle the case where an MFMA is copied to an AGPR, but not the case where the MFMA is copying from an AGPR.

perlfu

LGTM

Maybe change "fma" to "MFMA" in PR/commit title?

arsenm added the backend:AMDGPU label Aug 11, 2025 — with Graphite App

arsenm mentioned this pull request Aug 11, 2025

AMDGPU: Handle rewriting non-tied MFMA to AGPR form #153015

Merged

arsenm requested review from jrbyrnes, perlfu, qcolombet and srpande August 11, 2025 14:40

This was referenced Aug 11, 2025

AMDGPU: Handle rewriting VGPR MFMAs with immediate src2 #153016

Merged

AMDGPU: Add some baseline test for mfma rewrite with subregister copies #153018

Merged

This was referenced Aug 11, 2025

AMDGPU: Handle rewriting VGPR MFMA to AGPR with subregister copies #153019

Merged

AMDGPU: Add baseline test for mfma rewrite with phi #153021

Merged

arsenm marked this pull request as ready for review August 11, 2025 14:40

arsenm force-pushed the users/arsenm/amdgpu/add-baseline-test-mfma-src-is-agpr branch from 5ae2e08 to b5bbb0d Compare August 12, 2025 01:30

arsenm force-pushed the users/arsenm/amdgpu/allow-subregister-copies-mfma-rewrite-candidates branch from 3956361 to 603f8df Compare August 12, 2025 01:30

arsenm force-pushed the users/arsenm/amdgpu/add-baseline-test-mfma-src-is-agpr branch from b5bbb0d to 9e68d80 Compare August 18, 2025 15:31

arsenm force-pushed the users/arsenm/amdgpu/allow-subregister-copies-mfma-rewrite-candidates branch from 603f8df to d4de0b2 Compare August 18, 2025 15:31

This was referenced Aug 19, 2025

AMDGPU: Add baseline test for unspilling VGPRs after MFMA rewrite #154322

Open

AMDGPU: Try to unspill VGPRs after rewriting MFMAs to AGPR form #154323

Open

perlfu reviewed Aug 20, 2025

View reviewed changes

llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir Outdated Show resolved Hide resolved

arsenm force-pushed the users/arsenm/amdgpu/add-baseline-test-mfma-src-is-agpr branch from 9e68d80 to 6f0530c Compare August 20, 2025 23:23

arsenm force-pushed the users/arsenm/amdgpu/allow-subregister-copies-mfma-rewrite-candidates branch 2 times, most recently from 2ad172e to 7619736 Compare August 21, 2025 00:11

arsenm force-pushed the users/arsenm/amdgpu/add-baseline-test-mfma-src-is-agpr branch from 6f0530c to 7e1a5ca Compare August 21, 2025 00:11

AMDGPU: Handle rewriting VGPR MFMA to AGPR with subregister copies

05eb7bb

This should address the case where the result isn't fully used, resulting in partial copy bundles from the MFMA result.

arsenm added 3 commits August 21, 2025 09:41

Early exit if getLargestLegalSuperClass can't change the class

92a0a65

AMDGPU: Add baseline test for vgpr fma with copied-from AGPR

6463b0e

We currently handle the case where an MFMA is copied to an AGPR, but not the case where the MFMA is copying from an AGPR.

Fix comment

9ecd6f8

arsenm force-pushed the users/arsenm/amdgpu/allow-subregister-copies-mfma-rewrite-candidates branch from 7619736 to 92a0a65 Compare August 21, 2025 00:42

arsenm force-pushed the users/arsenm/amdgpu/add-baseline-test-mfma-src-is-agpr branch from 7e1a5ca to 9ecd6f8 Compare August 21, 2025 00:42

perlfu approved these changes Aug 21, 2025

View reviewed changes

arsenm changed the title ~~AMDGPU: Add baseline test for vgpr fma with copied-from AGPR~~ AMDGPU: Add baseline test for vgpr mfma with copied-from AGPR Aug 21, 2025

Base automatically changed from users/arsenm/amdgpu/allow-subregister-copies-mfma-rewrite-candidates to main August 21, 2025 01:17

arsenm merged commit bcf41e0 into main Aug 21, 2025
12 of 15 checks passed

arsenm deleted the users/arsenm/amdgpu/add-baseline-test-mfma-src-is-agpr branch August 21, 2025 01:24

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

AMDGPU: Add baseline test for vgpr mfma with copied-from AGPR #153020

AMDGPU: Add baseline test for vgpr mfma with copied-from AGPR #153020

Uh oh!

arsenm commented Aug 11, 2025

Uh oh!

arsenm commented Aug 11, 2025 •

edited

Loading

Uh oh!

llvmbot commented Aug 11, 2025

Uh oh!

Uh oh!

perlfu left a comment

Uh oh!

Uh oh!

Uh oh!

AMDGPU: Add baseline test for vgpr mfma with copied-from AGPR #153020

AMDGPU: Add baseline test for vgpr mfma with copied-from AGPR #153020

Uh oh!

Conversation

arsenm commented Aug 11, 2025

Uh oh!

arsenm commented Aug 11, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Aug 11, 2025

Uh oh!

Uh oh!

perlfu left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

arsenm commented Aug 11, 2025 •

edited

Loading