[llvm] [AMDGPU] Fixed verifier crash because of multiple live range components. (PR #190719)
Dhruva Chakrabarti via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 8 12:25:48 PDT 2026
https://github.com/dhruvachak updated https://github.com/llvm/llvm-project/pull/190719
>From 120df3edc94512782128accde914de6f933cc42d Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Fri, 3 Apr 2026 19:14:43 -0500
Subject: [PATCH 1/7] [AMDGPU] Fixed verifier crash because of multiple live
range components.
In Rewrite AGPR-Copy-MFMA pass, after replacing spill instructions, the
replacement register may have multiple live range components when the
spill slot was stored to more than once. The verifier crashes with a
bad machine code error. This patch fixes the problem by splitting a live
range but assigning the same physical register in this scenario. A new
test has been added that verifies the absence of this verifier error.
Assisted-by: Claude Opus
---
.../AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp | 12 ++
...ite-vgpr-mfma-to-agpr-spill-multi-store.ll | 146 ++++++++++++++++++
2 files changed, 158 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
index 7a5db42c7a89a..d8503cbab29e0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
@@ -548,7 +548,19 @@ void AMDGPURewriteAGPRCopyMFMAImpl::eliminateSpillsOfReassignedVGPRs() const {
// replacement vreg uses.
LiveInterval &NewLI = LIS.createAndComputeVirtRegInterval(NewVReg);
VRM.grow();
+
+ // A spill slot can be stored to multiple times, so the replacement
+ // vreg may have multiple disconnected live range components. Split
+ // them into separate vregs to maintain the single-component invariant.
+ SmallVector<LiveInterval *, 4> SplitLIs;
+ LIS.splitSeparateComponents(NewLI, SplitLIs);
+
LRM.assign(NewLI, PhysReg);
+ for (LiveInterval *SplitLI : SplitLIs) {
+ VRM.grow();
+ LRM.assign(*SplitLI, PhysReg);
+ }
+
MFI.RemoveStackObject(Slot);
break;
}
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
new file mode 100644
index 0000000000000..d6d9088033d15
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
@@ -0,0 +1,146 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -O3 \
+; RUN: -amdgpu-use-amdgpu-trackers=1 --lsr-drop-solution=1 \
+; RUN: -enable-post-misched=0 -verify-machineinstrs -o - %s 2>&1 | FileCheck %s
+; CHECK-NOT: Bad machine code: Multiple connected components in live interval
+
+; This test verifies that multiple connected live range components are not
+; created by the VGPR-to-AGPR MFMA rewrite pass. If multiple components exist,
+; the verifier would error out.
+
+define amdgpu_kernel void @multi_store_spill_slot() #0 {
+entry:
+ br label %do.body
+
+do.body:
+ %c_block_tile.sroa.37.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.70.0, %do.body ]
+ %c_block_tile.sroa.70.0 = phi float [ 0.000000e+00, %entry ], [ %ext0, %do.body ]
+ %c_block_tile.sroa.961.0 = phi float [ 0.000000e+00, %entry ], [ %ext1, %do.body ]
+ %c_block_tile.sroa.994.0 = phi float [ 0.000000e+00, %entry ], [ %ext2, %do.body ]
+ %c_block_tile.sroa.1060.0 = phi float [ 0.000000e+00, %entry ], [ %ext3, %do.body ]
+ %c_block_tile.sroa.1093.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.1060.0, %do.body ]
+ %c_block_tile.sroa.1588.0 = phi float [ 0.000000e+00, %entry ], [ %ext4, %do.body ]
+ %c_block_tile.sroa.1687.0 = phi float [ 0.000000e+00, %entry ], [ %ext5, %do.body ]
+ %c_block_tile.sroa.2578.0 = phi float [ 0.000000e+00, %entry ], [ %ext6, %do.body ]
+ %c_block_tile.sroa.2611.0 = phi float [ 0.000000e+00, %entry ], [ %ext7, %do.body ]
+ %c_block_tile.sroa.2644.0 = phi float [ 0.000000e+00, %entry ], [ %ext8, %do.body ]
+ %c_block_tile.sroa.2677.0 = phi float [ 0.000000e+00, %entry ], [ %ext9, %do.body ]
+ %c_block_tile.sroa.3568.0 = phi float [ 0.000000e+00, %entry ], [ %ext10, %do.body ]
+ %c_block_tile.sroa.3634.0 = phi float [ 0.000000e+00, %entry ], [ %ext11, %do.body ]
+ %c_block_tile.sroa.3898.0 = phi float [ 0.000000e+00, %entry ], [ %ext12, %do.body ]
+ %c_block_tile.sroa.3931.0 = phi float [ 0.000000e+00, %entry ], [ %ext13, %do.body ]
+ %c_block_tile.sroa.4624.0 = phi float [ 0.000000e+00, %entry ], [ %ext14, %do.body ]
+ %c_block_tile.sroa.4657.0 = phi float [ 0.000000e+00, %entry ], [ %ext15, %do.body ]
+ %c_block_tile.sroa.5185.0 = phi float [ 0.000000e+00, %entry ], [ %ext16, %do.body ]
+ %c_block_tile.sroa.5218.0 = phi float [ 0.000000e+00, %entry ], [ %ext17, %do.body ]
+ %c_block_tile.sroa.5746.0 = phi float [ 0.000000e+00, %entry ], [ %ext18, %do.body ]
+ %c_block_tile.sroa.5779.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.5746.0, %do.body ]
+ %c_block_tile.sroa.5812.0 = phi float [ 0.000000e+00, %entry ], [ %ext19, %do.body ]
+ %c_block_tile.sroa.6373.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.6406.0, %do.body ]
+ %c_block_tile.sroa.6406.0 = phi float [ 0.000000e+00, %entry ], [ %ext20, %do.body ]
+ %c_block_tile.sroa.7297.0 = phi float [ 0.000000e+00, %entry ], [ %ext21, %do.body ]
+ %c_block_tile.sroa.7363.0 = phi float [ 0.000000e+00, %entry ], [ %ext22, %do.body ]
+ %c_block_tile.sroa.7396.0 = phi float [ 0.000000e+00, %entry ], [ %ext23, %do.body ]
+ %c_block_tile.sroa.7429.0 = phi float [ 0.000000e+00, %entry ], [ %ext24, %do.body ]
+ %v0 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.37.0, i64 0
+ %v1 = insertelement <16 x float> %v0, float %c_block_tile.sroa.70.0, i64 0
+ %0 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v1, i32 0, i32 0, i32 0)
+ %v2 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.961.0, i64 13
+ %v3 = insertelement <16 x float> %v2, float %c_block_tile.sroa.994.0, i64 14
+ %v4 = insertelement <16 x float> %v3, float 0x7FF8000000000000, i64 0
+ %1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v4, i32 0, i32 0, i32 0)
+ %v5 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.1588.0, i64 0
+ %v6 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.1687.0, i64 3
+ %v7 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2578.0, i64 0
+ %v8 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2611.0, i64 0
+ %v9 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2677.0, i64 0
+ %v10 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.3568.0, i64 0
+ %v11 = insertelement <16 x float> splat (float 1.000000e+00), float %c_block_tile.sroa.3634.0, i64 0
+ %2 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v11, i32 0, i32 0, i32 0)
+ %v12 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.3898.0, i64 1
+ %v13 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.4624.0, i64 0
+ %v14 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5185.0, i64 0
+ %v15 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5218.0, i64 0
+ %v16 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5746.0, i64 14
+ %v17 = insertelement <16 x float> %v16, float %c_block_tile.sroa.5779.0, i64 15
+ %3 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v17, i32 0, i32 0, i32 0)
+ %v18 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5812.0, i64 0
+ %v19 = insertelement <16 x float> %v18, float %c_block_tile.sroa.5812.0, i64 0
+ %v20 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.6373.0, i64 1
+ %v21 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.7297.0, i64 0
+ %v22 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.7363.0, i64 15
+ %v23 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.7396.0, i64 0
+ %v24 = insertelement <16 x float> %v23, float %c_block_tile.sroa.7429.0, i64 1
+ %4 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v15, i32 0, i32 0, i32 0)
+ %5 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %4, i32 0, i32 0, i32 0)
+ %6 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v22, i32 0, i32 0, i32 0)
+ %7 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %6, i32 0, i32 0, i32 0)
+ %8 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v6, i32 0, i32 0, i32 0)
+ %9 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %0, i32 0, i32 0, i32 0)
+ %10 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %2, i32 0, i32 0, i32 0)
+ %v25 = insertelement <16 x float> %v12, float %c_block_tile.sroa.3931.0, i64 0
+ %11 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v25, i32 0, i32 0, i32 0)
+ %v26 = insertelement <16 x float> %v20, float %c_block_tile.sroa.6406.0, i64 0
+ %12 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v26, i32 0, i32 0, i32 0)
+ %13 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v24, i32 0, i32 0, i32 0)
+ %14 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %9, i32 0, i32 0, i32 0)
+ %15 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %11, i32 0, i32 0, i32 0)
+ %16 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %3, i32 0, i32 0, i32 0)
+ %17 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %1, i32 0, i32 0, i32 0)
+ %18 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %17, i32 0, i32 0, i32 0)
+ %v27 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.1060.0, i64 0
+ %v28 = insertelement <16 x float> %v27, float %c_block_tile.sroa.1093.0, i64 1
+ %19 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v28, i32 0, i32 0, i32 0)
+ %20 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %8, i32 0, i32 0, i32 0)
+ %21 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %20, i32 0, i32 0, i32 0)
+ %22 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %16, i32 0, i32 0, i32 0)
+ %23 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %12, i32 0, i32 0, i32 0)
+ %ext0 = extractelement <16 x float> %14, i64 0
+ %24 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %18, i32 0, i32 0, i32 0)
+ %ext1 = extractelement <16 x float> %24, i64 0
+ %ext2 = extractelement <16 x float> %17, i64 0
+ %25 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %19, i32 0, i32 0, i32 0)
+ %ext3 = extractelement <16 x float> %25, i64 0
+ %ext4 = extractelement <16 x float> %8, i64 0
+ %ext5 = extractelement <16 x float> %21, i64 0
+ %26 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v8, i32 0, i32 0, i32 0)
+ %ext6 = extractelement <16 x float> %26, i64 0
+ %27 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %26, i32 0, i32 0, i32 0)
+ %ext7 = extractelement <16 x float> %27, i64 0
+ %v29 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2644.0, i64 0
+ %28 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v29, i32 0, i32 0, i32 0)
+ %29 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %28, i32 0, i32 0, i32 0)
+ %ext8 = extractelement <16 x float> %29, i64 0
+ %ext9 = extractelement <16 x float> %28, i64 0
+ %ext10 = extractelement <16 x float> %2, i64 0
+ %30 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %10, i32 0, i32 0, i32 0)
+ %ext11 = extractelement <16 x float> %30, i64 0
+ %ext12 = extractelement <16 x float> %11, i64 0
+ %31 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %15, i32 0, i32 0, i32 0)
+ %32 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %31, i32 0, i32 0, i32 0)
+ %33 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> splat (half 0xH3C00), <16 x float> %32, i32 0, i32 0, i32 0)
+ %ext13 = extractelement <16 x float> %33, i64 7
+ %v30 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.4657.0, i64 0
+ %34 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v30, i32 0, i32 0, i32 0)
+ %ext14 = extractelement <16 x float> %34, i64 0
+ %35 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %34, i32 0, i32 0, i32 0)
+ %ext15 = extractelement <16 x float> %35, i64 0
+ %ext16 = extractelement <16 x float> %4, i64 0
+ %ext17 = extractelement <16 x float> %5, i64 0
+ %ext18 = extractelement <16 x float> %22, i64 0
+ %36 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v19, i32 0, i32 0, i32 0)
+ %ext19 = extractelement <16 x float> %36, i64 0
+ %37 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %23, i32 0, i32 0, i32 0)
+ %ext20 = extractelement <16 x float> %37, i64 0
+ %ext21 = extractelement <16 x float> %6, i64 0
+ %ext22 = extractelement <16 x float> %7, i64 0
+ %ext23 = extractelement <16 x float> %13, i64 0
+ %38 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %13, i32 0, i32 0, i32 0)
+ %ext24 = extractelement <16 x float> %38, i64 0
+ br label %do.body
+}
+
+declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half>, <4 x half>, <16 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+uselistorder ptr @llvm.amdgcn.mfma.f32.32x32x8f16, { 38, 37, 36, 35, 34, 33, 32, 31, 30, 29, 28, 27, 26, 25, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 13, 12, 11, 10, 9, 8, 7, 14, 6, 5, 4, 3, 2, 1, 0 }
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" }
>From 33ed5e893c3ea6bc7eb6d097d93dafebc5e819a8 Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Tue, 7 Apr 2026 11:31:25 -0500
Subject: [PATCH 2/7] [AMDGPU] Removed uselistorder from test, not required for
failure mode.
---
.../AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll | 2 --
1 file changed, 2 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
index d6d9088033d15..c2e9f335b02dc 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
@@ -141,6 +141,4 @@ do.body:
declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half>, <4 x half>, <16 x float>, i32 immarg, i32 immarg, i32 immarg)
-uselistorder ptr @llvm.amdgcn.mfma.f32.32x32x8f16, { 38, 37, 36, 35, 34, 33, 32, 31, 30, 29, 28, 27, 26, 25, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 13, 12, 11, 10, 9, 8, 7, 14, 6, 5, 4, 3, 2, 1, 0 }
-
attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" }
>From e6c262b5394c39cd89272bc89578f43243177c09 Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Tue, 7 Apr 2026 13:40:27 -0500
Subject: [PATCH 3/7] [AMDGPU] Added debugging output/test for multiple store
to spill slot.
---
llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp | 4 ++++
.../rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll | 10 +++++++---
2 files changed, 11 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
index d8503cbab29e0..6a03e9cbc66d3 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
@@ -555,6 +555,10 @@ void AMDGPURewriteAGPRCopyMFMAImpl::eliminateSpillsOfReassignedVGPRs() const {
SmallVector<LiveInterval *, 4> SplitLIs;
LIS.splitSeparateComponents(NewLI, SplitLIs);
+ LLVM_DEBUG(if (!SplitLIs.empty()) dbgs()
+ << "Split unspilled interval into " << (SplitLIs.size() + 1)
+ << " components\n");
+
LRM.assign(NewLI, PhysReg);
for (LiveInterval *SplitLI : SplitLIs) {
VRM.grow();
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
index c2e9f335b02dc..9250fb166b4fc 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
@@ -1,11 +1,15 @@
+; REQUIRES: asserts
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -O3 \
; RUN: -amdgpu-use-amdgpu-trackers=1 --lsr-drop-solution=1 \
-; RUN: -enable-post-misched=0 -verify-machineinstrs -o - %s 2>&1 | FileCheck %s
-; CHECK-NOT: Bad machine code: Multiple connected components in live interval
+; RUN: -enable-post-misched=0 -verify-machineinstrs \
+; RUN: -debug-only=amdgpu-rewrite-agpr-copy-mfma %s 2>&1 | FileCheck %s
; This test verifies that multiple connected live range components are not
; created by the VGPR-to-AGPR MFMA rewrite pass. If multiple components exist,
-; the verifier would error out.
+; the verifier would error out. Check that the unspilled interval was split
+; into separate components.
+
+; CHECK: Split unspilled interval into {{[0-9]+}} components
define amdgpu_kernel void @multi_store_spill_slot() #0 {
entry:
>From d4ed2a757b14bc5cfb21bac5aaea4e4e3dbc594f Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Tue, 7 Apr 2026 14:04:18 -0500
Subject: [PATCH 4/7] Trimmed test options and passes.
---
.../AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
index 9250fb166b4fc..6db1c56c1102f 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
@@ -1,7 +1,7 @@
; REQUIRES: asserts
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -O3 \
-; RUN: -amdgpu-use-amdgpu-trackers=1 --lsr-drop-solution=1 \
-; RUN: -enable-post-misched=0 -verify-machineinstrs \
+; RUN: -amdgpu-use-amdgpu-trackers=1 -verify-machineinstrs \
+; RUN: -stop-after=amdgpu-rewrite-agpr-copy-mfma \
; RUN: -debug-only=amdgpu-rewrite-agpr-copy-mfma %s 2>&1 | FileCheck %s
; This test verifies that multiple connected live range components are not
>From 7b16d100421189ea8585e46b85ed5e392a9e3659 Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Wed, 8 Apr 2026 10:47:09 -0500
Subject: [PATCH 5/7] Added braces.
---
llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp | 7 ++++---
1 file changed, 4 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
index 6a03e9cbc66d3..101986c34daec 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
@@ -555,9 +555,10 @@ void AMDGPURewriteAGPRCopyMFMAImpl::eliminateSpillsOfReassignedVGPRs() const {
SmallVector<LiveInterval *, 4> SplitLIs;
LIS.splitSeparateComponents(NewLI, SplitLIs);
- LLVM_DEBUG(if (!SplitLIs.empty()) dbgs()
- << "Split unspilled interval into " << (SplitLIs.size() + 1)
- << " components\n");
+ LLVM_DEBUG(if (!SplitLIs.empty()) {
+ dbgs() << "Split unspilled interval into " << (SplitLIs.size() + 1)
+ << " components\n";
+ });
LRM.assign(NewLI, PhysReg);
for (LiveInterval *SplitLI : SplitLIs) {
>From 8a699bfe2c14d81ebc280a130845b8e5c037a68c Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Wed, 8 Apr 2026 13:48:46 -0500
Subject: [PATCH 6/7] Added an MIR test with pruned phases.
---
...gpr-mfma-to-agpr-spill-multi-store-mir.mir | 772 ++++++++++++++++++
1 file changed, 772 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir
new file mode 100644
index 0000000000000..dd342fe241d9a
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir
@@ -0,0 +1,772 @@
+# REQUIRES: asserts
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+# RUN: -amdgpu-use-amdgpu-trackers=1 -verify-machineinstrs \
+# RUN: -start-before=register-coalescer \
+# RUN: -stop-after=amdgpu-rewrite-agpr-copy-mfma \
+# RUN: -debug-only=amdgpu-rewrite-agpr-copy-mfma -o /dev/null %s 2>&1 \
+# RUN: | FileCheck %s
+
+# This test verifies that when the VGPR-to-AGPR MFMA rewrite pass eliminates
+# spills of reassigned VGPRs, multiple connected live range components are
+# split into separate virtual registers to satisfy the verifier.
+
+# CHECK: Split unspilled interval into {{[0-9]+}} components
+
+--- |
+ define amdgpu_kernel void @multi_store_spill_slot() #0 {
+ entry:
+ br label %do.body
+
+ do.body:
+ br label %DummyReturnBlock
+
+ DummyReturnBlock:
+ unreachable
+ }
+
+ attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" "target-cpu"="gfx90a" }
+...
+---
+name: multi_store_spill_slot
+alignment: 4
+exposesReturnsTwice: false
+legalized: false
+regBankSelected: false
+selected: false
+failedISel: false
+tracksRegLiveness: true
+hasWinCFI: false
+noPhis: true
+isSSA: false
+noVRegs: false
+hasFakeUses: false
+callsEHReturn: false
+callsUnwindInit: false
+hasEHContTarget: false
+hasEHScopes: false
+hasEHFunclets: false
+isOutlined: false
+debugInstrRef: false
+failsVerification: false
+tracksDebugUserValues: false
+registers:
+ - { id: 0, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 1, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 2, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 3, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 4, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 5, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 6, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 7, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 8, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 9, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 10, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 11, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 12, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 13, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 14, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 15, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 16, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 17, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 18, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 19, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 20, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 21, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 22, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 23, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 24, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 25, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 26, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 27, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 28, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 29, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 30, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 31, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 32, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 33, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 34, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 35, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 36, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 37, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 38, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 39, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 40, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 41, class: sgpr_128, preferred-register: '', flags: [ ] }
+ - { id: 42, class: sgpr_64, preferred-register: '', flags: [ ] }
+ - { id: 43, class: sgpr_64, preferred-register: '', flags: [ ] }
+ - { id: 44, class: sgpr_64, preferred-register: '', flags: [ ] }
+ - { id: 45, class: sgpr_64, preferred-register: '', flags: [ ] }
+ - { id: 46, class: sgpr_64, preferred-register: '', flags: [ ] }
+ - { id: 47, class: sgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 48, class: sgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 49, class: sgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 50, class: sgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 51, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 52, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 53, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 54, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 55, class: sreg_32, preferred-register: '', flags: [ ] }
+ - { id: 56, class: sreg_64, preferred-register: '', flags: [ ] }
+ - { id: 57, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 58, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 59, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 60, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 61, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 62, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 63, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 64, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 65, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 66, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 67, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 68, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 69, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 70, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 71, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 72, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 73, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 74, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 75, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 76, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 77, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 78, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 79, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 80, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 81, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 82, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 83, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 84, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 85, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 86, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 87, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 88, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 89, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 90, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 91, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 92, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 93, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 94, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 95, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 96, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 97, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 98, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 99, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 100, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 101, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 102, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 103, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 104, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 105, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 106, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 107, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 108, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 109, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 110, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 111, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 112, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 113, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 114, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 115, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 116, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 117, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 118, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 119, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 120, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 121, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 122, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 123, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 124, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 125, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 126, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 127, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 128, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 129, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 130, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 131, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 132, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 133, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 134, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 135, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 136, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 137, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 138, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 139, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 140, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 141, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 142, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 143, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 144, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 145, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 146, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 147, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 148, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 149, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 150, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 151, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 152, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 153, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 154, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 155, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 156, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 157, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 158, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 159, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 160, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 161, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 162, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 163, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 164, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 165, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 166, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 167, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 168, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 169, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 170, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 171, class: sreg_32, preferred-register: '', flags: [ ] }
+ - { id: 172, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 173, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 174, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 175, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 176, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 177, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 178, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 179, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 180, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 181, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 182, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 183, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 184, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 185, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 186, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 187, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 188, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 189, class: vreg_512_align2, preferred-register: '', flags: [ ] }
+ - { id: 190, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 191, class: av_64_align2, preferred-register: '', flags: [ ] }
+ - { id: 192, class: sreg_64, preferred-register: '', flags: [ ] }
+ - { id: 193, class: sreg_64, preferred-register: '', flags: [ ] }
+ - { id: 194, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 195, class: vgpr_32, preferred-register: '', flags: [ ] }
+ - { id: 196, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 197, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 198, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 199, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 200, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 201, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 202, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 203, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 204, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 205, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 206, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 207, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 208, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 209, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 210, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 211, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 212, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 213, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 214, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 215, class: av_32, preferred-register: '', flags: [ ] }
+ - { id: 216, class: av_32, preferred-register: '', flags: [ ] }
+liveins: []
+frameInfo:
+ isFrameAddressTaken: false
+ isReturnAddressTaken: false
+ hasStackMap: false
+ hasPatchPoint: false
+ stackSize: 0
+ offsetAdjustment: 0
+ maxAlignment: 1
+ adjustsStack: false
+ hasCalls: false
+ stackProtector: ''
+ functionContext: ''
+ maxCallFrameSize: 4294967295
+ cvBytesOfCalleeSavedRegisters: 0
+ hasOpaqueSPAdjustment: false
+ hasVAStart: false
+ hasMustTailInVarArgFunc: false
+ hasTailCall: false
+ isCalleeSavedInfoValid: false
+ localFrameSize: 0
+fixedStack: []
+stack: []
+entry_values: []
+callSites: []
+debugValueSubstitutions: []
+constants: []
+machineFunctionInfo:
+ explicitKernArgSize: 0
+ maxKernArgAlign: 8
+ ldsSize: 32768
+ gdsSize: 0
+ dynLDSAlign: 1
+ isEntryFunction: true
+ isChainFunction: false
+ memoryBound: false
+ waveLimiter: false
+ hasSpilledSGPRs: false
+ hasSpilledVGPRs: false
+ numWaveDispatchSGPRs: 0
+ numWaveDispatchVGPRs: 0
+ scratchRSrcReg: '$sgpr96_sgpr97_sgpr98_sgpr99'
+ frameOffsetReg: '$fp_reg'
+ stackPtrOffsetReg: '$sgpr32'
+ bytesInStackArgArea: 0
+ returnsVoid: true
+ argumentInfo:
+ privateSegmentBuffer: { reg: '$sgpr0_sgpr1_sgpr2_sgpr3' }
+ dispatchPtr: { reg: '$sgpr4_sgpr5' }
+ queuePtr: { reg: '$sgpr6_sgpr7' }
+ kernargSegmentPtr: { reg: '$sgpr8_sgpr9' }
+ dispatchID: { reg: '$sgpr10_sgpr11' }
+ flatScratchInit: { reg: '$sgpr12_sgpr13' }
+ workGroupIDX: { reg: '$sgpr14' }
+ workGroupIDY: { reg: '$sgpr15' }
+ workGroupIDZ: { reg: '$sgpr16' }
+ privateSegmentWaveByteOffset: { reg: '$sgpr17' }
+ workItemIDX: { reg: '$vgpr0', mask: 1023 }
+ workItemIDY: { reg: '$vgpr0', mask: 1047552 }
+ workItemIDZ: { reg: '$vgpr0', mask: 1072693248 }
+ psInputAddr: 0
+ psInputEnable: 0
+ maxMemoryClusterDWords: 8
+ mode:
+ ieee: true
+ dx10-clamp: true
+ fp32-input-denormals: true
+ fp32-output-denormals: true
+ fp64-fp16-input-denormals: true
+ fp64-fp16-output-denormals: true
+ highBitsOf32BitAddress: 0
+ occupancy: 2
+ vgprForAGPRCopy: ''
+ sgprForEXECCopy: '$sgpr100_sgpr101'
+ longBranchReservedReg: ''
+ hasInitWholeWave: false
+ dynamicVGPRBlockSize: 0
+ scratchReservedForDynamicVGPRs: 0
+ numKernargPreloadSGPRs: 0
+ isWholeWaveFunction: false
+body: |
+ bb.0.entry:
+ successors: %bb.1(0x80000000)
+
+ %52:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %55:sreg_32 = S_MOV_B32 0
+ undef %56.sub0:sreg_64 = COPY %55
+ %56.sub1:sreg_64 = COPY killed %55
+ %58:av_64_align2 = COPY killed %56
+ %60:vgpr_32 = V_MOV_B32_e32 2143289344, implicit $exec
+ %67:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 1065353216, implicit $exec
+ %194:vgpr_32 = V_MOV_B32_e32 1006648320, implicit $exec
+ undef %175.sub0:av_64_align2 = COPY %194
+ %175.sub1:av_64_align2 = COPY killed %194
+ %193:sreg_64 = S_AND_B64 $exec, -1, implicit-def dead $scc
+ %196:av_32 = COPY %52
+ %197:av_32 = COPY %52
+ %198:av_32 = COPY %52
+ %199:av_32 = COPY %52
+ %200:av_32 = COPY %52
+ %201:av_32 = COPY %52
+ %202:av_32 = COPY %52
+ %203:av_32 = COPY %52
+ %204:av_32 = COPY %52
+ %205:av_32 = COPY %52
+ %206:av_32 = COPY %52
+ %207:av_32 = COPY %52
+ %208:av_32 = COPY %52
+ %209:av_32 = COPY %52
+ %210:av_32 = COPY %52
+ %211:av_32 = COPY %52
+ %212:av_32 = COPY %52
+ %213:av_32 = COPY %52
+ %214:av_32 = COPY %52
+ %215:av_32 = COPY %52
+ %216:av_32 = COPY %52
+
+ bb.1.do.body:
+ successors: %bb.2(0x04000000), %bb.1(0x7c000000)
+
+ %20:av_32 = COPY killed %215
+ %19:av_32 = COPY killed %214
+ %18:av_32 = COPY killed %213
+ %17:av_32 = COPY killed %212
+ %16:av_32 = COPY killed %211
+ %15:av_32 = COPY killed %210
+ %14:av_32 = COPY killed %209
+ %13:av_32 = COPY killed %208
+ %12:av_32 = COPY killed %207
+ %11:av_32 = COPY killed %206
+ %10:av_32 = COPY killed %205
+ %9:av_32 = COPY killed %204
+ %8:av_32 = COPY killed %203
+ %7:av_32 = COPY killed %202
+ %6:av_32 = COPY killed %201
+ %5:av_32 = COPY killed %200
+ %4:av_32 = COPY killed %199
+ %3:av_32 = COPY killed %198
+ %2:av_32 = COPY killed %197
+ %1:av_32 = COPY killed %196
+ undef %54.sub0:vreg_512_align2 = COPY killed %1
+ %54.sub1:vreg_512_align2 = COPY %52
+ %54.sub2:vreg_512_align2 = COPY %52
+ %54.sub3:vreg_512_align2 = COPY %52
+ %54.sub4:vreg_512_align2 = COPY %52
+ %54.sub5:vreg_512_align2 = COPY %52
+ %54.sub6:vreg_512_align2 = COPY %52
+ %54.sub7:vreg_512_align2 = COPY %52
+ %54.sub8:vreg_512_align2 = COPY %52
+ %54.sub9:vreg_512_align2 = COPY %52
+ %54.sub10:vreg_512_align2 = COPY %52
+ %54.sub11:vreg_512_align2 = COPY %52
+ %54.sub12:vreg_512_align2 = COPY %52
+ %54.sub13:vreg_512_align2 = COPY %52
+ %54.sub14:vreg_512_align2 = COPY %52
+ %54.sub15:vreg_512_align2 = COPY %52
+ %57:vreg_512_align2 = COPY killed %54
+ %57:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %57, 0, 0, 0, implicit $mode, implicit $exec
+ undef %61.sub0:vreg_512_align2 = COPY %60
+ %61.sub1:vreg_512_align2 = COPY %52
+ %61.sub2:vreg_512_align2 = COPY %52
+ %61.sub3:vreg_512_align2 = COPY %52
+ %61.sub4:vreg_512_align2 = COPY %52
+ %61.sub5:vreg_512_align2 = COPY %52
+ %61.sub6:vreg_512_align2 = COPY %52
+ %61.sub7:vreg_512_align2 = COPY %52
+ %61.sub8:vreg_512_align2 = COPY %52
+ %61.sub9:vreg_512_align2 = COPY %52
+ %61.sub10:vreg_512_align2 = COPY %52
+ %61.sub11:vreg_512_align2 = COPY %52
+ %61.sub12:vreg_512_align2 = COPY %52
+ %61.sub13:vreg_512_align2 = COPY killed %2
+ %61.sub14:vreg_512_align2 = COPY killed %3
+ %61.sub15:vreg_512_align2 = COPY %52
+ %62:vreg_512_align2 = COPY killed %61
+ %62:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %62, 0, 0, 0, implicit $mode, implicit $exec
+ undef %65.sub0:vreg_512_align2 = COPY %52
+ %65.sub1:vreg_512_align2 = COPY %52
+ %65.sub2:vreg_512_align2 = COPY %52
+ %65.sub3:vreg_512_align2 = COPY killed %6
+ %65.sub4:vreg_512_align2 = COPY %52
+ %65.sub5:vreg_512_align2 = COPY %52
+ %65.sub6:vreg_512_align2 = COPY %52
+ %65.sub7:vreg_512_align2 = COPY %52
+ %65.sub8:vreg_512_align2 = COPY %52
+ %65.sub9:vreg_512_align2 = COPY %52
+ %65.sub10:vreg_512_align2 = COPY %52
+ %65.sub11:vreg_512_align2 = COPY %52
+ %65.sub12:vreg_512_align2 = COPY %52
+ %65.sub13:vreg_512_align2 = COPY %52
+ %65.sub14:vreg_512_align2 = COPY %52
+ %65.sub15:vreg_512_align2 = COPY %52
+ undef %66.sub0:vreg_512_align2 = COPY killed %7
+ %66.sub1:vreg_512_align2 = COPY %52
+ %66.sub2:vreg_512_align2 = COPY %52
+ %66.sub3:vreg_512_align2 = COPY %52
+ %66.sub4:vreg_512_align2 = COPY %52
+ %66.sub5:vreg_512_align2 = COPY %52
+ %66.sub6:vreg_512_align2 = COPY %52
+ %66.sub7:vreg_512_align2 = COPY %52
+ %66.sub8:vreg_512_align2 = COPY %52
+ %66.sub9:vreg_512_align2 = COPY %52
+ %66.sub10:vreg_512_align2 = COPY %52
+ %66.sub11:vreg_512_align2 = COPY %52
+ %66.sub12:vreg_512_align2 = COPY %52
+ %66.sub13:vreg_512_align2 = COPY %52
+ %66.sub14:vreg_512_align2 = COPY %52
+ %66.sub15:vreg_512_align2 = COPY %52
+ undef %68.sub0:vreg_512_align2 = COPY killed %9
+ %68.sub1:vreg_512_align2 = COPY %67
+ %68.sub2:vreg_512_align2 = COPY %67
+ %68.sub3:vreg_512_align2 = COPY %67
+ %68.sub4:vreg_512_align2 = COPY %67
+ %68.sub5:vreg_512_align2 = COPY %67
+ %68.sub6:vreg_512_align2 = COPY %67
+ %68.sub7:vreg_512_align2 = COPY %67
+ %68.sub8:vreg_512_align2 = COPY %67
+ %68.sub9:vreg_512_align2 = COPY %67
+ %68.sub10:vreg_512_align2 = COPY %67
+ %68.sub11:vreg_512_align2 = COPY %67
+ %68.sub12:vreg_512_align2 = COPY %67
+ %68.sub13:vreg_512_align2 = COPY %67
+ %68.sub14:vreg_512_align2 = COPY %67
+ %68.sub15:vreg_512_align2 = COPY %67
+ %69:vreg_512_align2 = COPY killed %68
+ %69:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %69, 0, 0, 0, implicit $mode, implicit $exec
+ undef %72.sub0:vreg_512_align2 = COPY killed %11
+ %72.sub1:vreg_512_align2 = COPY killed %10
+ %72.sub2:vreg_512_align2 = COPY %52
+ %72.sub3:vreg_512_align2 = COPY %52
+ %72.sub4:vreg_512_align2 = COPY %52
+ %72.sub5:vreg_512_align2 = COPY %52
+ %72.sub6:vreg_512_align2 = COPY %52
+ %72.sub7:vreg_512_align2 = COPY %52
+ %72.sub8:vreg_512_align2 = COPY %52
+ %72.sub9:vreg_512_align2 = COPY %52
+ %72.sub10:vreg_512_align2 = COPY %52
+ %72.sub11:vreg_512_align2 = COPY %52
+ %72.sub12:vreg_512_align2 = COPY %52
+ %72.sub13:vreg_512_align2 = COPY %52
+ %72.sub14:vreg_512_align2 = COPY %52
+ %72.sub15:vreg_512_align2 = COPY %52
+ undef %73.sub0:vreg_512_align2 = COPY killed %13
+ %73.sub1:vreg_512_align2 = COPY %52
+ %73.sub2:vreg_512_align2 = COPY %52
+ %73.sub3:vreg_512_align2 = COPY %52
+ %73.sub4:vreg_512_align2 = COPY %52
+ %73.sub5:vreg_512_align2 = COPY %52
+ %73.sub6:vreg_512_align2 = COPY %52
+ %73.sub7:vreg_512_align2 = COPY %52
+ %73.sub8:vreg_512_align2 = COPY %52
+ %73.sub9:vreg_512_align2 = COPY %52
+ %73.sub10:vreg_512_align2 = COPY %52
+ %73.sub11:vreg_512_align2 = COPY %52
+ %73.sub12:vreg_512_align2 = COPY %52
+ %73.sub13:vreg_512_align2 = COPY %52
+ %73.sub14:vreg_512_align2 = COPY %52
+ %73.sub15:vreg_512_align2 = COPY %52
+ undef %74.sub0:vreg_512_align2 = COPY %52
+ %74.sub1:vreg_512_align2 = COPY %52
+ %74.sub2:vreg_512_align2 = COPY %52
+ %74.sub3:vreg_512_align2 = COPY %52
+ %74.sub4:vreg_512_align2 = COPY %52
+ %74.sub5:vreg_512_align2 = COPY %52
+ %74.sub6:vreg_512_align2 = COPY %52
+ %74.sub7:vreg_512_align2 = COPY %52
+ %74.sub8:vreg_512_align2 = COPY %52
+ %74.sub9:vreg_512_align2 = COPY %52
+ %74.sub10:vreg_512_align2 = COPY %52
+ %74.sub11:vreg_512_align2 = COPY %52
+ %74.sub12:vreg_512_align2 = COPY %52
+ %74.sub13:vreg_512_align2 = COPY %52
+ %74.sub14:vreg_512_align2 = COPY %14
+ %74.sub15:vreg_512_align2 = COPY killed %15
+ %75:vreg_512_align2 = COPY killed %74
+ %75:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %75, 0, 0, 0, implicit $mode, implicit $exec
+ undef %78.sub0:vreg_512_align2 = COPY killed %16
+ %78.sub1:vreg_512_align2 = COPY %52
+ %78.sub2:vreg_512_align2 = COPY %52
+ %78.sub3:vreg_512_align2 = COPY %52
+ %78.sub4:vreg_512_align2 = COPY %52
+ %78.sub5:vreg_512_align2 = COPY %52
+ %78.sub6:vreg_512_align2 = COPY %52
+ %78.sub7:vreg_512_align2 = COPY %52
+ %78.sub8:vreg_512_align2 = COPY %52
+ %78.sub9:vreg_512_align2 = COPY %52
+ %78.sub10:vreg_512_align2 = COPY %52
+ %78.sub11:vreg_512_align2 = COPY %52
+ %78.sub12:vreg_512_align2 = COPY %52
+ %78.sub13:vreg_512_align2 = COPY %52
+ %78.sub14:vreg_512_align2 = COPY %52
+ %78.sub15:vreg_512_align2 = COPY %52
+ undef %79.sub0:vreg_512_align2 = COPY %18
+ %79.sub1:vreg_512_align2 = COPY killed %17
+ %79.sub2:vreg_512_align2 = COPY %52
+ %79.sub3:vreg_512_align2 = COPY %52
+ %79.sub4:vreg_512_align2 = COPY %52
+ %79.sub5:vreg_512_align2 = COPY %52
+ %79.sub6:vreg_512_align2 = COPY %52
+ %79.sub7:vreg_512_align2 = COPY %52
+ %79.sub8:vreg_512_align2 = COPY %52
+ %79.sub9:vreg_512_align2 = COPY %52
+ %79.sub10:vreg_512_align2 = COPY %52
+ %79.sub11:vreg_512_align2 = COPY %52
+ %79.sub12:vreg_512_align2 = COPY %52
+ %79.sub13:vreg_512_align2 = COPY %52
+ %79.sub14:vreg_512_align2 = COPY %52
+ %79.sub15:vreg_512_align2 = COPY %52
+ undef %80.sub0:vreg_512_align2 = COPY %52
+ %80.sub1:vreg_512_align2 = COPY %52
+ %80.sub2:vreg_512_align2 = COPY %52
+ %80.sub3:vreg_512_align2 = COPY %52
+ %80.sub4:vreg_512_align2 = COPY %52
+ %80.sub5:vreg_512_align2 = COPY %52
+ %80.sub6:vreg_512_align2 = COPY %52
+ %80.sub7:vreg_512_align2 = COPY %52
+ %80.sub8:vreg_512_align2 = COPY %52
+ %80.sub9:vreg_512_align2 = COPY %52
+ %80.sub10:vreg_512_align2 = COPY %52
+ %80.sub11:vreg_512_align2 = COPY %52
+ %80.sub12:vreg_512_align2 = COPY %52
+ %80.sub13:vreg_512_align2 = COPY %52
+ %80.sub14:vreg_512_align2 = COPY %52
+ %80.sub15:vreg_512_align2 = COPY killed %19
+ %21:av_32 = COPY killed %216
+ undef %81.sub0:vreg_512_align2 = COPY killed %20
+ %81.sub1:vreg_512_align2 = COPY killed %21
+ %81.sub2:vreg_512_align2 = COPY %52
+ %81.sub3:vreg_512_align2 = COPY %52
+ %81.sub4:vreg_512_align2 = COPY %52
+ %81.sub5:vreg_512_align2 = COPY %52
+ %81.sub6:vreg_512_align2 = COPY %52
+ %81.sub7:vreg_512_align2 = COPY %52
+ %81.sub8:vreg_512_align2 = COPY %52
+ %81.sub9:vreg_512_align2 = COPY %52
+ %81.sub10:vreg_512_align2 = COPY %52
+ %81.sub11:vreg_512_align2 = COPY %52
+ %81.sub12:vreg_512_align2 = COPY %52
+ %81.sub13:vreg_512_align2 = COPY %52
+ %81.sub14:vreg_512_align2 = COPY %52
+ %81.sub15:vreg_512_align2 = COPY %52
+ %82:vreg_512_align2 = COPY killed %73
+ %82:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %82, 0, 0, 0, implicit $mode, implicit $exec
+ %85:vreg_512_align2 = COPY killed %82
+ %85:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %85, 0, 0, 0, implicit $mode, implicit $exec
+ %88:vreg_512_align2 = COPY killed %80
+ %88:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %88, 0, 0, 0, implicit $mode, implicit $exec
+ %91:vreg_512_align2 = COPY killed %88
+ %91:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %91, 0, 0, 0, implicit $mode, implicit $exec
+ %94:vreg_512_align2 = COPY killed %65
+ %94:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %94, 0, 0, 0, implicit $mode, implicit $exec
+ %97:vreg_512_align2 = COPY killed %57
+ %97:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %97, 0, 0, 0, implicit $mode, implicit $exec
+ %100:vreg_512_align2 = COPY killed %69
+ %100:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %100, 0, 0, 0, implicit $mode, implicit $exec
+ %103:vreg_512_align2 = COPY killed %72
+ %103:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %103, 0, 0, 0, implicit $mode, implicit $exec
+ %106:vreg_512_align2 = COPY killed %79
+ %106:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %106, 0, 0, 0, implicit $mode, implicit $exec
+ %109:vreg_512_align2 = COPY killed %81
+ %109:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %109, 0, 0, 0, implicit $mode, implicit $exec
+ %112:vreg_512_align2 = COPY killed %97
+ %112:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %112, 0, 0, 0, implicit $mode, implicit $exec
+ early-clobber %115:vreg_512_align2 = V_MFMA_F32_32X32X8F16_vgprcd_e64 %58, %58, %103, 0, 0, 0, implicit $mode, implicit $exec
+ %118:vreg_512_align2 = COPY killed %75
+ %118:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %118, 0, 0, 0, implicit $mode, implicit $exec
+ %121:vreg_512_align2 = COPY killed %62
+ %121:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %121, 0, 0, 0, implicit $mode, implicit $exec
+ early-clobber %124:vreg_512_align2 = V_MFMA_F32_32X32X8F16_vgprcd_e64 %58, %58, %121, 0, 0, 0, implicit $mode, implicit $exec
+ undef %127.sub0:vreg_512_align2 = COPY %4
+ %127.sub1:vreg_512_align2 = COPY killed %5
+ %127.sub2:vreg_512_align2 = COPY %52
+ %127.sub3:vreg_512_align2 = COPY %52
+ %127.sub4:vreg_512_align2 = COPY %52
+ %127.sub5:vreg_512_align2 = COPY %52
+ %127.sub6:vreg_512_align2 = COPY %52
+ %127.sub7:vreg_512_align2 = COPY %52
+ %127.sub8:vreg_512_align2 = COPY %52
+ %127.sub9:vreg_512_align2 = COPY %52
+ %127.sub10:vreg_512_align2 = COPY %52
+ %127.sub11:vreg_512_align2 = COPY %52
+ %127.sub12:vreg_512_align2 = COPY %52
+ %127.sub13:vreg_512_align2 = COPY %52
+ %127.sub14:vreg_512_align2 = COPY %52
+ %127.sub15:vreg_512_align2 = COPY %52
+ %128:vreg_512_align2 = COPY killed %127
+ %128:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %128, 0, 0, 0, implicit $mode, implicit $exec
+ %131:vreg_512_align2 = COPY killed %94
+ %131:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %131, 0, 0, 0, implicit $mode, implicit $exec
+ %134:vreg_512_align2 = COPY killed %131
+ %134:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %134, 0, 0, 0, implicit $mode, implicit $exec
+ %137:vreg_512_align2 = COPY killed %118
+ %137:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %137, 0, 0, 0, implicit $mode, implicit $exec
+ %140:vreg_512_align2 = COPY killed %106
+ %140:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %140, 0, 0, 0, implicit $mode, implicit $exec
+ %22:av_32 = COPY killed %112.sub0
+ %143:vreg_512_align2 = COPY killed %124
+ %143:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %143, 0, 0, 0, implicit $mode, implicit $exec
+ %23:av_32 = COPY killed %143.sub0
+ %24:av_32 = COPY killed %121.sub0
+ %146:vreg_512_align2 = COPY killed %128
+ %146:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %146, 0, 0, 0, implicit $mode, implicit $exec
+ %25:av_32 = COPY killed %146.sub0
+ %26:av_32 = COPY killed %134.sub0
+ %149:vreg_512_align2 = COPY killed %66
+ %149:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %149, 0, 0, 0, implicit $mode, implicit $exec
+ %152:vreg_512_align2 = COPY killed %149
+ %152:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %152, 0, 0, 0, implicit $mode, implicit $exec
+ %27:av_32 = COPY killed %152.sub0
+ undef %155.sub0:vreg_512_align2 = COPY killed %8
+ %155.sub1:vreg_512_align2 = COPY %52
+ %155.sub2:vreg_512_align2 = COPY %52
+ %155.sub3:vreg_512_align2 = COPY %52
+ %155.sub4:vreg_512_align2 = COPY %52
+ %155.sub5:vreg_512_align2 = COPY %52
+ %155.sub6:vreg_512_align2 = COPY %52
+ %155.sub7:vreg_512_align2 = COPY %52
+ %155.sub8:vreg_512_align2 = COPY %52
+ %155.sub9:vreg_512_align2 = COPY %52
+ %155.sub10:vreg_512_align2 = COPY %52
+ %155.sub11:vreg_512_align2 = COPY %52
+ %155.sub12:vreg_512_align2 = COPY %52
+ %155.sub13:vreg_512_align2 = COPY %52
+ %155.sub14:vreg_512_align2 = COPY %52
+ %155.sub15:vreg_512_align2 = COPY %52
+ %156:vreg_512_align2 = COPY killed %155
+ %156:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %156, 0, 0, 0, implicit $mode, implicit $exec
+ %159:vreg_512_align2 = COPY killed %156
+ %159:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %159, 0, 0, 0, implicit $mode, implicit $exec
+ %28:av_32 = COPY killed %159.sub0
+ %162:vreg_512_align2 = COPY killed %100
+ %162:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %162, 0, 0, 0, implicit $mode, implicit $exec
+ %29:av_32 = COPY killed %162.sub0
+ %30:av_32 = COPY killed %103.sub0
+ %165:vreg_512_align2 = COPY killed %115
+ %165:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %165, 0, 0, 0, implicit $mode, implicit $exec
+ %168:vreg_512_align2 = COPY killed %165
+ %168:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %168, 0, 0, 0, implicit $mode, implicit $exec
+ %173:vreg_512_align2 = COPY killed %168
+ %173:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %175, %173, 0, 0, 0, implicit $mode, implicit $exec
+ %31:av_32 = COPY killed %173.sub7
+ undef %176.sub0:vreg_512_align2 = COPY killed %12
+ %176.sub1:vreg_512_align2 = COPY %52
+ %176.sub2:vreg_512_align2 = COPY %52
+ %176.sub3:vreg_512_align2 = COPY %52
+ %176.sub4:vreg_512_align2 = COPY %52
+ %176.sub5:vreg_512_align2 = COPY %52
+ %176.sub6:vreg_512_align2 = COPY %52
+ %176.sub7:vreg_512_align2 = COPY %52
+ %176.sub8:vreg_512_align2 = COPY %52
+ %176.sub9:vreg_512_align2 = COPY %52
+ %176.sub10:vreg_512_align2 = COPY %52
+ %176.sub11:vreg_512_align2 = COPY %52
+ %176.sub12:vreg_512_align2 = COPY %52
+ %176.sub13:vreg_512_align2 = COPY %52
+ %176.sub14:vreg_512_align2 = COPY %52
+ %176.sub15:vreg_512_align2 = COPY %52
+ %177:vreg_512_align2 = COPY killed %176
+ %177:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %177, 0, 0, 0, implicit $mode, implicit $exec
+ %180:vreg_512_align2 = COPY killed %177
+ %180:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %180, 0, 0, 0, implicit $mode, implicit $exec
+ %32:av_32 = COPY killed %180.sub0
+ %33:av_32 = COPY killed %85.sub0
+ %34:av_32 = COPY killed %137.sub0
+ %183:vreg_512_align2 = COPY killed %78
+ %183:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %183, 0, 0, 0, implicit $mode, implicit $exec
+ %35:av_32 = COPY killed %183.sub0
+ %186:vreg_512_align2 = COPY killed %140
+ %186:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %186, 0, 0, 0, implicit $mode, implicit $exec
+ %36:av_32 = COPY killed %186.sub0
+ %37:av_32 = COPY killed %91.sub0
+ %38:av_32 = COPY %109.sub0
+ %189:vreg_512_align2 = COPY killed %109
+ %189:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %58, %58, %189, 0, 0, 0, implicit $mode, implicit $exec
+ %39:av_32 = COPY killed %189.sub0
+ $vcc = COPY %193
+ %196:av_32 = COPY killed %22
+ %197:av_32 = COPY killed %23
+ %198:av_32 = COPY killed %24
+ %199:av_32 = COPY killed %25
+ %200:av_32 = COPY killed %4
+ %201:av_32 = COPY killed %26
+ %202:av_32 = COPY killed %27
+ %203:av_32 = COPY killed %28
+ %204:av_32 = COPY killed %29
+ %205:av_32 = COPY killed %30
+ %206:av_32 = COPY killed %31
+ %207:av_32 = COPY killed %32
+ %208:av_32 = COPY killed %33
+ %209:av_32 = COPY killed %34
+ %210:av_32 = COPY killed %14
+ %211:av_32 = COPY killed %35
+ %212:av_32 = COPY killed %18
+ %213:av_32 = COPY killed %36
+ %214:av_32 = COPY killed %37
+ %215:av_32 = COPY killed %38
+ %216:av_32 = COPY killed %39
+ S_CBRANCH_VCCNZ %bb.1, implicit killed $vcc
+ S_BRANCH %bb.2
+
+ bb.2.DummyReturnBlock:
+ S_ENDPGM 0
+...
>From 33f20f6e89f1a36864c76746f1f2fd7bdbb79faf Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Wed, 8 Apr 2026 14:22:58 -0500
Subject: [PATCH 7/7] Added braces for better formatting.
---
llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp | 8 +++++---
1 file changed, 5 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
index 101986c34daec..5e0c845b2515b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
@@ -555,9 +555,11 @@ void AMDGPURewriteAGPRCopyMFMAImpl::eliminateSpillsOfReassignedVGPRs() const {
SmallVector<LiveInterval *, 4> SplitLIs;
LIS.splitSeparateComponents(NewLI, SplitLIs);
- LLVM_DEBUG(if (!SplitLIs.empty()) {
- dbgs() << "Split unspilled interval into " << (SplitLIs.size() + 1)
- << " components\n";
+ LLVM_DEBUG({
+ if (!SplitLIs.empty()) {
+ dbgs() << "Split unspilled interval into " << (SplitLIs.size() + 1)
+ << " components\n";
+ }
});
LRM.assign(NewLI, PhysReg);
More information about the llvm-commits
mailing list