Revert "[AMDGPU] Fixed verifier crash because of multiple live range components." (#193135)

Reverts llvm/llvm-project#190719

The Buildbot has detected a new failure on builder
sanitizer-aarch64-linux-bootstrap-hwasan while building llvm.

3 files changed, 0 insertions, 626 deletions

diff --git a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
index 1d39b4f1bc52..53be369db35c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURewriteAGPRCopyMFMA.cpp
@@ -554,26 +554,7 @@ 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);
-
-      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();
-        LRM.assign(*SplitLI, PhysReg);
-      }
-
       MFI.RemoveStackObject(Slot);
       break;
     }
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
deleted file mode 100644
index 1309dea96b27..000000000000
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir
+++ /dev/null
@@ -1,459 +0,0 @@
-# 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 -filetype=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
-tracksRegLiveness: true
-noPhis:          true
-isSSA:           false
-machineFunctionInfo:
-  ldsSize:         32768
-  stackPtrOffsetReg: '$sgpr32'
-body:             |
-  bb.0.entry:
-    successors: %bb.1(0x80000000)
-  
-    %39:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
-    %41:sreg_32 = S_MOV_B32 0
-    undef %42.sub0:sreg_64 = COPY %41
-    %42.sub1:sreg_64 = COPY killed %41
-    %44:av_64_align2 = COPY killed %42
-    %45:vgpr_32 = V_MOV_B32_e32 2143289344, implicit $exec
-    %50:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 1065353216, implicit $exec
-    %101:vgpr_32 = V_MOV_B32_e32 1006648320, implicit $exec
-    undef %93.sub0:av_64_align2 = COPY %101
-    %93.sub1:av_64_align2 = COPY killed %101
-    %100:sreg_64 = S_AND_B64 $exec, -1, implicit-def dead $scc
-    %102:av_32 = COPY %39
-    %103:av_32 = COPY %39
-    %104:av_32 = COPY %39
-    %105:av_32 = COPY %39
-    %106:av_32 = COPY %39
-    %107:av_32 = COPY %39
-    %108:av_32 = COPY %39
-    %109:av_32 = COPY %39
-    %110:av_32 = COPY %39
-    %111:av_32 = COPY %39
-    %112:av_32 = COPY %39
-    %113:av_32 = COPY %39
-    %114:av_32 = COPY %39
-    %115:av_32 = COPY %39
-    %116:av_32 = COPY %39
-    %117:av_32 = COPY %39
-    %118:av_32 = COPY %39
-    %119:av_32 = COPY %39
-    %120:av_32 = COPY %39
-    %121:av_32 = COPY %39
-    %122:av_32 = COPY %39
-  
-  bb.1.do.body:
-    successors: %bb.2(0x04000000), %bb.1(0x7c000000)
-  
-    %19:av_32 = COPY killed %121
-    %18:av_32 = COPY killed %120
-    %17:av_32 = COPY killed %119
-    %16:av_32 = COPY killed %118
-    %15:av_32 = COPY killed %117
-    %14:av_32 = COPY killed %116
-    %13:av_32 = COPY killed %115
-    %12:av_32 = COPY killed %114
-    %11:av_32 = COPY killed %113
-    %10:av_32 = COPY killed %112
-    %9:av_32 = COPY killed %111
-    %8:av_32 = COPY killed %110
-    %7:av_32 = COPY killed %109
-    %6:av_32 = COPY killed %108
-    %5:av_32 = COPY killed %107
-    %4:av_32 = COPY killed %106
-    %3:av_32 = COPY killed %105
-    %2:av_32 = COPY killed %104
-    %1:av_32 = COPY killed %103
-    %0:av_32 = COPY killed %102
-    undef %40.sub0:vreg_512_align2 = COPY killed %0
-    %40.sub1:vreg_512_align2 = COPY %39
-    %40.sub2:vreg_512_align2 = COPY %39
-    %40.sub3:vreg_512_align2 = COPY %39
-    %40.sub4:vreg_512_align2 = COPY %39
-    %40.sub5:vreg_512_align2 = COPY %39
-    %40.sub6:vreg_512_align2 = COPY %39
-    %40.sub7:vreg_512_align2 = COPY %39
-    %40.sub8:vreg_512_align2 = COPY %39
-    %40.sub9:vreg_512_align2 = COPY %39
-    %40.sub10:vreg_512_align2 = COPY %39
-    %40.sub11:vreg_512_align2 = COPY %39
-    %40.sub12:vreg_512_align2 = COPY %39
-    %40.sub13:vreg_512_align2 = COPY %39
-    %40.sub14:vreg_512_align2 = COPY %39
-    %40.sub15:vreg_512_align2 = COPY %39
-    %43:vreg_512_align2 = COPY killed %40
-    %43:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %43, 0, 0, 0, implicit $mode, implicit $exec
-    undef %46.sub0:vreg_512_align2 = COPY %45
-    %46.sub1:vreg_512_align2 = COPY %39
-    %46.sub2:vreg_512_align2 = COPY %39
-    %46.sub3:vreg_512_align2 = COPY %39
-    %46.sub4:vreg_512_align2 = COPY %39
-    %46.sub5:vreg_512_align2 = COPY %39
-    %46.sub6:vreg_512_align2 = COPY %39
-    %46.sub7:vreg_512_align2 = COPY %39
-    %46.sub8:vreg_512_align2 = COPY %39
-    %46.sub9:vreg_512_align2 = COPY %39
-    %46.sub10:vreg_512_align2 = COPY %39
-    %46.sub11:vreg_512_align2 = COPY %39
-    %46.sub12:vreg_512_align2 = COPY %39
-    %46.sub13:vreg_512_align2 = COPY killed %1
-    %46.sub14:vreg_512_align2 = COPY killed %2
-    %46.sub15:vreg_512_align2 = COPY %39
-    %47:vreg_512_align2 = COPY killed %46
-    %47:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %47, 0, 0, 0, implicit $mode, implicit $exec
-    undef %48.sub0:vreg_512_align2 = COPY %39
-    %48.sub1:vreg_512_align2 = COPY %39
-    %48.sub2:vreg_512_align2 = COPY %39
-    %48.sub3:vreg_512_align2 = COPY killed %5
-    %48.sub4:vreg_512_align2 = COPY %39
-    %48.sub5:vreg_512_align2 = COPY %39
-    %48.sub6:vreg_512_align2 = COPY %39
-    %48.sub7:vreg_512_align2 = COPY %39
-    %48.sub8:vreg_512_align2 = COPY %39
-    %48.sub9:vreg_512_align2 = COPY %39
-    %48.sub10:vreg_512_align2 = COPY %39
-    %48.sub11:vreg_512_align2 = COPY %39
-    %48.sub12:vreg_512_align2 = COPY %39
-    %48.sub13:vreg_512_align2 = COPY %39
-    %48.sub14:vreg_512_align2 = COPY %39
-    %48.sub15:vreg_512_align2 = COPY %39
-    undef %49.sub0:vreg_512_align2 = COPY killed %6
-    %49.sub1:vreg_512_align2 = COPY %39
-    %49.sub2:vreg_512_align2 = COPY %39
-    %49.sub3:vreg_512_align2 = COPY %39
-    %49.sub4:vreg_512_align2 = COPY %39
-    %49.sub5:vreg_512_align2 = COPY %39
-    %49.sub6:vreg_512_align2 = COPY %39
-    %49.sub7:vreg_512_align2 = COPY %39
-    %49.sub8:vreg_512_align2 = COPY %39
-    %49.sub9:vreg_512_align2 = COPY %39
-    %49.sub10:vreg_512_align2 = COPY %39
-    %49.sub11:vreg_512_align2 = COPY %39
-    %49.sub12:vreg_512_align2 = COPY %39
-    %49.sub13:vreg_512_align2 = COPY %39
-    %49.sub14:vreg_512_align2 = COPY %39
-    %49.sub15:vreg_512_align2 = COPY %39
-    undef %51.sub0:vreg_512_align2 = COPY killed %8
-    %51.sub1:vreg_512_align2 = COPY %50
-    %51.sub2:vreg_512_align2 = COPY %50
-    %51.sub3:vreg_512_align2 = COPY %50
-    %51.sub4:vreg_512_align2 = COPY %50
-    %51.sub5:vreg_512_align2 = COPY %50
-    %51.sub6:vreg_512_align2 = COPY %50
-    %51.sub7:vreg_512_align2 = COPY %50
-    %51.sub8:vreg_512_align2 = COPY %50
-    %51.sub9:vreg_512_align2 = COPY %50
-    %51.sub10:vreg_512_align2 = COPY %50
-    %51.sub11:vreg_512_align2 = COPY %50
-    %51.sub12:vreg_512_align2 = COPY %50
-    %51.sub13:vreg_512_align2 = COPY %50
-    %51.sub14:vreg_512_align2 = COPY %50
-    %51.sub15:vreg_512_align2 = COPY %50
-    %52:vreg_512_align2 = COPY killed %51
-    %52:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %52, 0, 0, 0, implicit $mode, implicit $exec
-    undef %53.sub0:vreg_512_align2 = COPY killed %10
-    %53.sub1:vreg_512_align2 = COPY killed %9
-    %53.sub2:vreg_512_align2 = COPY %39
-    %53.sub3:vreg_512_align2 = COPY %39
-    %53.sub4:vreg_512_align2 = COPY %39
-    %53.sub5:vreg_512_align2 = COPY %39
-    %53.sub6:vreg_512_align2 = COPY %39
-    %53.sub7:vreg_512_align2 = COPY %39
-    %53.sub8:vreg_512_align2 = COPY %39
-    %53.sub9:vreg_512_align2 = COPY %39
-    %53.sub10:vreg_512_align2 = COPY %39
-    %53.sub11:vreg_512_align2 = COPY %39
-    %53.sub12:vreg_512_align2 = COPY %39
-    %53.sub13:vreg_512_align2 = COPY %39
-    %53.sub14:vreg_512_align2 = COPY %39
-    %53.sub15:vreg_512_align2 = COPY %39
-    undef %54.sub0:vreg_512_align2 = COPY killed %12
-    %54.sub1:vreg_512_align2 = COPY %39
-    %54.sub2:vreg_512_align2 = COPY %39
-    %54.sub3:vreg_512_align2 = COPY %39
-    %54.sub4:vreg_512_align2 = COPY %39
-    %54.sub5:vreg_512_align2 = COPY %39
-    %54.sub6:vreg_512_align2 = COPY %39
-    %54.sub7:vreg_512_align2 = COPY %39
-    %54.sub8:vreg_512_align2 = COPY %39
-    %54.sub9:vreg_512_align2 = COPY %39
-    %54.sub10:vreg_512_align2 = COPY %39
-    %54.sub11:vreg_512_align2 = COPY %39
-    %54.sub12:vreg_512_align2 = COPY %39
-    %54.sub13:vreg_512_align2 = COPY %39
-    %54.sub14:vreg_512_align2 = COPY %39
-    %54.sub15:vreg_512_align2 = COPY %39
-    undef %55.sub0:vreg_512_align2 = COPY %39
-    %55.sub1:vreg_512_align2 = COPY %39
-    %55.sub2:vreg_512_align2 = COPY %39
-    %55.sub3:vreg_512_align2 = COPY %39
-    %55.sub4:vreg_512_align2 = COPY %39
-    %55.sub5:vreg_512_align2 = COPY %39
-    %55.sub6:vreg_512_align2 = COPY %39
-    %55.sub7:vreg_512_align2 = COPY %39
-    %55.sub8:vreg_512_align2 = COPY %39
-    %55.sub9:vreg_512_align2 = COPY %39
-    %55.sub10:vreg_512_align2 = COPY %39
-    %55.sub11:vreg_512_align2 = COPY %39
-    %55.sub12:vreg_512_align2 = COPY %39
-    %55.sub13:vreg_512_align2 = COPY %39
-    %55.sub14:vreg_512_align2 = COPY %13
-    %55.sub15:vreg_512_align2 = COPY killed %14
-    %56:vreg_512_align2 = COPY killed %55
-    %56:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %56, 0, 0, 0, implicit $mode, implicit $exec
-    undef %57.sub0:vreg_512_align2 = COPY killed %15
-    %57.sub1:vreg_512_align2 = COPY %39
-    %57.sub2:vreg_512_align2 = COPY %39
-    %57.sub3:vreg_512_align2 = COPY %39
-    %57.sub4:vreg_512_align2 = COPY %39
-    %57.sub5:vreg_512_align2 = COPY %39
-    %57.sub6:vreg_512_align2 = COPY %39
-    %57.sub7:vreg_512_align2 = COPY %39
-    %57.sub8:vreg_512_align2 = COPY %39
-    %57.sub9:vreg_512_align2 = COPY %39
-    %57.sub10:vreg_512_align2 = COPY %39
-    %57.sub11:vreg_512_align2 = COPY %39
-    %57.sub12:vreg_512_align2 = COPY %39
-    %57.sub13:vreg_512_align2 = COPY %39
-    %57.sub14:vreg_512_align2 = COPY %39
-    %57.sub15:vreg_512_align2 = COPY %39
-    undef %58.sub0:vreg_512_align2 = COPY %17
-    %58.sub1:vreg_512_align2 = COPY killed %16
-    %58.sub2:vreg_512_align2 = COPY %39
-    %58.sub3:vreg_512_align2 = COPY %39
-    %58.sub4:vreg_512_align2 = COPY %39
-    %58.sub5:vreg_512_align2 = COPY %39
-    %58.sub6:vreg_512_align2 = COPY %39
-    %58.sub7:vreg_512_align2 = COPY %39
-    %58.sub8:vreg_512_align2 = COPY %39
-    %58.sub9:vreg_512_align2 = COPY %39
-    %58.sub10:vreg_512_align2 = COPY %39
-    %58.sub11:vreg_512_align2 = COPY %39
-    %58.sub12:vreg_512_align2 = COPY %39
-    %58.sub13:vreg_512_align2 = COPY %39
-    %58.sub14:vreg_512_align2 = COPY %39
-    %58.sub15:vreg_512_align2 = COPY %39
-    undef %59.sub0:vreg_512_align2 = COPY %39
-    %59.sub1:vreg_512_align2 = COPY %39
-    %59.sub2:vreg_512_align2 = COPY %39
-    %59.sub3:vreg_512_align2 = COPY %39
-    %59.sub4:vreg_512_align2 = COPY %39
-    %59.sub5:vreg_512_align2 = COPY %39
-    %59.sub6:vreg_512_align2 = COPY %39
-    %59.sub7:vreg_512_align2 = COPY %39
-    %59.sub8:vreg_512_align2 = COPY %39
-    %59.sub9:vreg_512_align2 = COPY %39
-    %59.sub10:vreg_512_align2 = COPY %39
-    %59.sub11:vreg_512_align2 = COPY %39
-    %59.sub12:vreg_512_align2 = COPY %39
-    %59.sub13:vreg_512_align2 = COPY %39
-    %59.sub14:vreg_512_align2 = COPY %39
-    %59.sub15:vreg_512_align2 = COPY killed %18
-    %20:av_32 = COPY killed %122
-    undef %60.sub0:vreg_512_align2 = COPY killed %19
-    %60.sub1:vreg_512_align2 = COPY killed %20
-    %60.sub2:vreg_512_align2 = COPY %39
-    %60.sub3:vreg_512_align2 = COPY %39
-    %60.sub4:vreg_512_align2 = COPY %39
-    %60.sub5:vreg_512_align2 = COPY %39
-    %60.sub6:vreg_512_align2 = COPY %39
-    %60.sub7:vreg_512_align2 = COPY %39
-    %60.sub8:vreg_512_align2 = COPY %39
-    %60.sub9:vreg_512_align2 = COPY %39
-    %60.sub10:vreg_512_align2 = COPY %39
-    %60.sub11:vreg_512_align2 = COPY %39
-    %60.sub12:vreg_512_align2 = COPY %39
-    %60.sub13:vreg_512_align2 = COPY %39
-    %60.sub14:vreg_512_align2 = COPY %39
-    %60.sub15:vreg_512_align2 = COPY %39
-    %61:vreg_512_align2 = COPY killed %54
-    %61:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %61, 0, 0, 0, implicit $mode, implicit $exec
-    %62:vreg_512_align2 = COPY killed %61
-    %62:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %62, 0, 0, 0, implicit $mode, implicit $exec
-    %63:vreg_512_align2 = COPY killed %59
-    %63:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %63, 0, 0, 0, implicit $mode, implicit $exec
-    %64:vreg_512_align2 = COPY killed %63
-    %64:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %64, 0, 0, 0, implicit $mode, implicit $exec
-    %65:vreg_512_align2 = COPY killed %48
-    %65:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %65, 0, 0, 0, implicit $mode, implicit $exec
-    %66:vreg_512_align2 = COPY killed %43
-    %66:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %66, 0, 0, 0, implicit $mode, implicit $exec
-    %67:vreg_512_align2 = COPY killed %52
-    %67:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %67, 0, 0, 0, implicit $mode, implicit $exec
-    %68:vreg_512_align2 = COPY killed %53
-    %68:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %68, 0, 0, 0, implicit $mode, implicit $exec
-    %69:vreg_512_align2 = COPY killed %58
-    %69:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %69, 0, 0, 0, implicit $mode, implicit $exec
-    %70:vreg_512_align2 = COPY killed %60
-    %70:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %70, 0, 0, 0, implicit $mode, implicit $exec
-    %71:vreg_512_align2 = COPY killed %66
-    %71:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %71, 0, 0, 0, implicit $mode, implicit $exec
-    early-clobber %72:vreg_512_align2 = V_MFMA_F32_32X32X8F16_vgprcd_e64 %44, %44, %68, 0, 0, 0, implicit $mode, implicit $exec
-    %73:vreg_512_align2 = COPY killed %56
-    %73:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %73, 0, 0, 0, implicit $mode, implicit $exec
-    %74:vreg_512_align2 = COPY killed %47
-    %74:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %74, 0, 0, 0, implicit $mode, implicit $exec
-    early-clobber %75:vreg_512_align2 = V_MFMA_F32_32X32X8F16_vgprcd_e64 %44, %44, %74, 0, 0, 0, implicit $mode, implicit $exec
-    undef %76.sub0:vreg_512_align2 = COPY %3
-    %76.sub1:vreg_512_align2 = COPY killed %4
-    %76.sub2:vreg_512_align2 = COPY %39
-    %76.sub3:vreg_512_align2 = COPY %39
-    %76.sub4:vreg_512_align2 = COPY %39
-    %76.sub5:vreg_512_align2 = COPY %39
-    %76.sub6:vreg_512_align2 = COPY %39
-    %76.sub7:vreg_512_align2 = COPY %39
-    %76.sub8:vreg_512_align2 = COPY %39
-    %76.sub9:vreg_512_align2 = COPY %39
-    %76.sub10:vreg_512_align2 = COPY %39
-    %76.sub11:vreg_512_align2 = COPY %39
-    %76.sub12:vreg_512_align2 = COPY %39
-    %76.sub13:vreg_512_align2 = COPY %39
-    %76.sub14:vreg_512_align2 = COPY %39
-    %76.sub15:vreg_512_align2 = COPY %39
-    %77:vreg_512_align2 = COPY killed %76
-    %77:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %77, 0, 0, 0, implicit $mode, implicit $exec
-    %78:vreg_512_align2 = COPY killed %65
-    %78:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %78, 0, 0, 0, implicit $mode, implicit $exec
-    %79:vreg_512_align2 = COPY killed %78
-    %79:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %79, 0, 0, 0, implicit $mode, implicit $exec
-    %80:vreg_512_align2 = COPY killed %73
-    %80:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %80, 0, 0, 0, implicit $mode, implicit $exec
-    %81:vreg_512_align2 = COPY killed %69
-    %81:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %81, 0, 0, 0, implicit $mode, implicit $exec
-    %21:av_32 = COPY killed %71.sub0
-    %82:vreg_512_align2 = COPY killed %75
-    %82:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %82, 0, 0, 0, implicit $mode, implicit $exec
-    %22:av_32 = COPY killed %82.sub0
-    %23:av_32 = COPY killed %74.sub0
-    %83:vreg_512_align2 = COPY killed %77
-    %83:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %83, 0, 0, 0, implicit $mode, implicit $exec
-    %24:av_32 = COPY killed %83.sub0
-    %25:av_32 = COPY killed %79.sub0
-    %84:vreg_512_align2 = COPY killed %49
-    %84:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %84, 0, 0, 0, implicit $mode, implicit $exec
-    %85:vreg_512_align2 = COPY killed %84
-    %85:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %85, 0, 0, 0, implicit $mode, implicit $exec
-    %26:av_32 = COPY killed %85.sub0
-    undef %86.sub0:vreg_512_align2 = COPY killed %7
-    %86.sub1:vreg_512_align2 = COPY %39
-    %86.sub2:vreg_512_align2 = COPY %39
-    %86.sub3:vreg_512_align2 = COPY %39
-    %86.sub4:vreg_512_align2 = COPY %39
-    %86.sub5:vreg_512_align2 = COPY %39
-    %86.sub6:vreg_512_align2 = COPY %39
-    %86.sub7:vreg_512_align2 = COPY %39
-    %86.sub8:vreg_512_align2 = COPY %39
-    %86.sub9:vreg_512_align2 = COPY %39
-    %86.sub10:vreg_512_align2 = COPY %39
-    %86.sub11:vreg_512_align2 = COPY %39
-    %86.sub12:vreg_512_align2 = COPY %39
-    %86.sub13:vreg_512_align2 = COPY %39
-    %86.sub14:vreg_512_align2 = COPY %39
-    %86.sub15:vreg_512_align2 = COPY %39
-    %87:vreg_512_align2 = COPY killed %86
-    %87:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %87, 0, 0, 0, implicit $mode, implicit $exec
-    %88:vreg_512_align2 = COPY killed %87
-    %88:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %88, 0, 0, 0, implicit $mode, implicit $exec
-    %27:av_32 = COPY killed %88.sub0
-    %89:vreg_512_align2 = COPY killed %67
-    %89:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %89, 0, 0, 0, implicit $mode, implicit $exec
-    %28:av_32 = COPY killed %89.sub0
-    %29:av_32 = COPY killed %68.sub0
-    %90:vreg_512_align2 = COPY killed %72
-    %90:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %90, 0, 0, 0, implicit $mode, implicit $exec
-    %91:vreg_512_align2 = COPY killed %90
-    %91:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %91, 0, 0, 0, implicit $mode, implicit $exec
-    %92:vreg_512_align2 = COPY killed %91
-    %92:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %93, %92, 0, 0, 0, implicit $mode, implicit $exec
-    %30:av_32 = COPY killed %92.sub7
-    undef %94.sub0:vreg_512_align2 = COPY killed %11
-    %94.sub1:vreg_512_align2 = COPY %39
-    %94.sub2:vreg_512_align2 = COPY %39
-    %94.sub3:vreg_512_align2 = COPY %39
-    %94.sub4:vreg_512_align2 = COPY %39
-    %94.sub5:vreg_512_align2 = COPY %39
-    %94.sub6:vreg_512_align2 = COPY %39
-    %94.sub7:vreg_512_align2 = COPY %39
-    %94.sub8:vreg_512_align2 = COPY %39
-    %94.sub9:vreg_512_align2 = COPY %39
-    %94.sub10:vreg_512_align2 = COPY %39
-    %94.sub11:vreg_512_align2 = COPY %39
-    %94.sub12:vreg_512_align2 = COPY %39
-    %94.sub13:vreg_512_align2 = COPY %39
-    %94.sub14:vreg_512_align2 = COPY %39
-    %94.sub15:vreg_512_align2 = COPY %39
-    %95:vreg_512_align2 = COPY killed %94
-    %95:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %95, 0, 0, 0, implicit $mode, implicit $exec
-    %96:vreg_512_align2 = COPY killed %95
-    %96:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %96, 0, 0, 0, implicit $mode, implicit $exec
-    %31:av_32 = COPY killed %96.sub0
-    %32:av_32 = COPY killed %62.sub0
-    %33:av_32 = COPY killed %80.sub0
-    %97:vreg_512_align2 = COPY killed %57
-    %97:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %97, 0, 0, 0, implicit $mode, implicit $exec
-    %34:av_32 = COPY killed %97.sub0
-    %98:vreg_512_align2 = COPY killed %81
-    %98:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %98, 0, 0, 0, implicit $mode, implicit $exec
-    %35:av_32 = COPY killed %98.sub0
-    %36:av_32 = COPY killed %64.sub0
-    %37:av_32 = COPY %70.sub0
-    %99:vreg_512_align2 = COPY killed %70
-    %99:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %99, 0, 0, 0, implicit $mode, implicit $exec
-    %38:av_32 = COPY killed %99.sub0
-    $vcc = COPY %100
-    %102:av_32 = COPY killed %21
-    %103:av_32 = COPY killed %22
-    %104:av_32 = COPY killed %23
-    %105:av_32 = COPY killed %24
-    %106:av_32 = COPY killed %3
-    %107:av_32 = COPY killed %25
-    %108:av_32 = COPY killed %26
-    %109:av_32 = COPY killed %27
-    %110:av_32 = COPY killed %28
-    %111:av_32 = COPY killed %29
-    %112:av_32 = COPY killed %30
-    %113:av_32 = COPY killed %31
-    %114:av_32 = COPY killed %32
-    %115:av_32 = COPY killed %33
-    %116:av_32 = COPY killed %13
-    %117:av_32 = COPY killed %34
-    %118:av_32 = COPY killed %17
-    %119:av_32 = COPY killed %35
-    %120:av_32 = COPY killed %36
-    %121:av_32 = COPY killed %37
-    %122:av_32 = COPY killed %38
-    S_CBRANCH_VCCNZ %bb.1, implicit killed $vcc
-    S_BRANCH %bb.2
-  
-  bb.2.DummyReturnBlock:
-    S_ENDPGM 0
-...
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
deleted file mode 100644
index 6db1c56c1102..000000000000
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
+++ /dev/null
@@ -1,148 +0,0 @@
-; REQUIRES: asserts
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -O3 \
-; 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
-; created by the VGPR-to-AGPR MFMA rewrite pass. If multiple components exist,
-; 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:
-  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)
-
-attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" }