[AMDGPU] Reduce AGPR to AGPR copies with same source

[vangthao]

Summary

On targets where direct AGPR to AGPR copy is not possible, the postrapseudos
pass tries to expand AGPR to AGPR copies by first trying to find the defining
accvgpr_write instruction and propagating its vgpr register to avoid creating
temporary registers. In cases where we cannot propagate, a new temporary VGPR
register will be created for each AGPR copy instruction even if the sources
are the same for all or some of the copy instructions.

With this patch, after failing to propagate the defining instruction's VGPR
register, we will next try to see if another copy instruction has already
inserted in a temporay VGPR copy with accvgpr_read and attempt to propagate
that copy.

[arsenm]

I think copyPhysReg should not have to do any nontrivial scanning. These situations where this matters should be resolvable before register allocation

[rampitec]

In D105062#2845855, @arsenm wrote:

I think copyPhysReg should not have to do any nontrivial scanning. These situations where this matters should be resolvable before register allocation

I concur. This is untrivial and error prone. There are quite a number of foldings for agpr copy and init in different places, this looks like a candidate for yet another one, starting from a relevant testcase. This place is really a last resort.

Do you have a minimal ir test showing a problem you are solving?

[vangthao]

In D105062#2845881, @rampitec wrote:

In D105062#2845855, @arsenm wrote:

I think copyPhysReg should not have to do any nontrivial scanning. These situations where this matters should be resolvable before register allocation

I concur. This is untrivial and error prone. There are quite a number of foldings for agpr copy and init in different places, this looks like a candidate for yet another one, starting from a relevant testcase. This place is really a last resort.

Do you have a minimal ir test showing a problem you are solving?

I added a test case broadcast_a32_unable_to_propagate_v0 in accvgpr-copy.mir where we failed to propagate v0 when broadcasting $agpr0 to 31 other lanes. Using the trunk, this forces us to read $agpr0 31 times, write it to a vgpr register, and then copy it to the agpr lane. This issue occurs when an instruction is scheduled in-between the defining accvgpr_write and the COPIES and also writes over the same vgpr register as the defining accvgpr_write. In this case, a V_MOV_B32_e32 uses $vgpr0 and thus we cannot propagate $vgpr0 to the AGPR copies.

I believe this is really only an issue when we have a situation where multiple lanes are using the same source. Following the IR printed by LLVM before register allocation, the code looks somewhat like this

%80:vgpr_32 = COPY %12:sgpr_32
%81:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
%82:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
%83:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
%84:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
%85:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
...
%70:areg_1024 = REG_SEQUENCE %81:agpr_32, %subreg.sub0, %82:agpr_32 ...

This seems fine to me. We are using the vgpr register instead of an agpr register before register allocation. But then the machine-cse pass eliminates this chain of instruction and creates a dependency on agpr register %81.

%80:vgpr_32 = COPY %12:sgpr_32
%81:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
%70:areg_1024 = REG_SEQUENCE %81:agpr_32, %subreg.sub0, %81:agpr_32, %subreg.sub1, %81:agpr_32 ...

This is an issue because the REG_SEQUENCE will be expanded into multiple AGPR copies during twoaddressinstruction pass.

%80:vgpr_32 = V_MOV_B32_e32 1123418112, implicit $exec
%81:agpr_32 = V_ACCVGPR_WRITE_B32_e64 killed %80:vgpr_32, implicit $exec
undef %70.sub0:areg_1024 = COPY %81:agpr_32
%70.sub1:areg_1024 = COPY %81:agpr_32
%70.sub2:areg_1024 = COPY %81:agpr_32
%70.sub3:areg_1024 = COPY %81:agpr_32
...

So we just opted to eliminate instructions and replaced them with AGPR copies that we will have to transform back into those same instructions we eliminated. Even worse if we happen to not be able to propagate %80:vgpr_32 as in the test scenario broadcast_a32_unable_to_propagate_v0.

Disabling machine-cse gives us something like this:

64B       %80:vgpr_32 = V_MOV_B32_e32 1123418112, implicit $exec
80B       %81:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
96B       %82:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
112B      %83:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
128B      %84:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
144B      %85:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
...
608B      %70.sub1:areg_1024 = COPY %82:agpr_32
624B      %70.sub2:areg_1024 = COPY %83:agpr_32
640B      %70.sub3:areg_1024 = COPY %84:agpr_32
656B      %70.sub4:areg_1024 = COPY %85:agpr_32
672B      %70.sub5:areg_1024 = COPY %86:agpr_32
...

Which then the AGPR copies is folded into the accvgpr_write by the simple register coalescing pass.

64B       %80:vgpr_32 = V_MOV_B32_e32 1123418112, implicit $exec
80B       undef %70.sub0:areg_1024 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
96B       %70.sub1:areg_1024 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
112B      %70.sub2:areg_1024 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
128B      %70.sub3:areg_1024 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
144B      %70.sub4:areg_1024 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
...

This seems to be perfectly fine and is what we wanted since there are no AGPR copies and the accvgpr_writes are using vgpr sources. So I believe the issue here is that the machine-cse pass is removing something that we wish to keep and replacing it with something possibly worse. Eliminating accvgpr_write and replacing them with AGPR copies does not seem to be equivalent or better on targets that are not able to copy from AGPR to AGPR directly.

[rampitec]

In D105062#2845965, @vangthao wrote:

This is an issue because the REG_SEQUENCE will be expanded into multiple AGPR copies during twoaddressinstruction pass.

%80:vgpr_32 = V_MOV_B32_e32 1123418112, implicit $exec
%81:agpr_32 = V_ACCVGPR_WRITE_B32_e64 killed %80:vgpr_32, implicit $exec
undef %70.sub0:areg_1024 = COPY %81:agpr_32
%70.sub1:areg_1024 = COPY %81:agpr_32
%70.sub2:areg_1024 = COPY %81:agpr_32
%70.sub3:areg_1024 = COPY %81:agpr_32
...

Can we somehow catch it at this point? This seems to be a culprit.

[vangthao]

In D105062#2847773, @rampitec wrote:

In D105062#2845965, @vangthao wrote:

This is an issue because the REG_SEQUENCE will be expanded into multiple AGPR copies during twoaddressinstruction pass.

%80:vgpr_32 = V_MOV_B32_e32 1123418112, implicit $exec
%81:agpr_32 = V_ACCVGPR_WRITE_B32_e64 killed %80:vgpr_32, implicit $exec
undef %70.sub0:areg_1024 = COPY %81:agpr_32
%70.sub1:areg_1024 = COPY %81:agpr_32
%70.sub2:areg_1024 = COPY %81:agpr_32
%70.sub3:areg_1024 = COPY %81:agpr_32
...

Can we somehow catch it at this point? This seems to be a culprit.

I agree this does seem to be the culprit so we would need to catch it between this pass and register allocation. However, I'm not sure which pass in-between those two should have the responsibility of catching this or whether there should be a new pass to handle this.

[rampitec]

In D105062#2850276, @vangthao wrote:

%80:vgpr_32 = COPY %12:sgpr_32
%81:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %80:vgpr_32, implicit $exec
%70:areg_1024 = REG_SEQUENCE %81:agpr_32, %subreg.sub0, %81:agpr_32, %subreg.sub1, %81:agpr_32 ...

I agree this does seem to be the culprit so we would need to catch it between this pass and register allocation. However, I'm not sure which pass in-between those two should have the responsibility of catching this or whether there should be a new pass to handle this.

It looks like a special case where we splat an sgpr into an areg... What will happen if we catch it in FixSGPRCopies and keep reg_sequence initializing areg directly from the sgpr?

The other possibility is to try catching it before two address pass in the SIFoldOperands and attempt to fold vgpr into reg_sequence. It should possible to play with mir tests to see how it will work.

[rampitec]

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

581

In any way this comment will need to be removed and function renamed if the problem is resolved.

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

static void indirectCopyToAGPR(const SIInstrInfo &TII,
                                RegScavenger &RS,
                                Register ImpDefSuperReg = Register(),
                                Register ImpUseSuperReg = Register()) {
   const SIRegisterInfo &RI = TII.getRegisterInfo();
 
   assert(AMDGPU::SReg_32RegClass.contains(SrcReg) ||
          AMDGPU::AGPR_32RegClass.contains(SrcReg));
 
+  MachineBasicBlock::iterator DefPoint = MBB.begin();
   // First try to find defining accvgpr_write to avoid temporary registers.
   for (auto Def = MI, E = MBB.begin(); Def != E; ) {
     --Def;
     if (!Def->definesRegister(SrcReg, &RI))
       continue;
     if (Def->getOpcode() != AMDGPU::V_ACCVGPR_WRITE_B32_e64)
       break;
 
     MachineOperand &DefOp = Def->getOperand(1);
     assert(DefOp.isReg() || DefOp.isImm());
 
     if (DefOp.isReg()) {
       // Check that register source operand if not clobbered before MI.
       // Immediate operands are always safe to propagate.
       bool SafeToPropagate = true;
       for (auto I = Def; I != MI && SafeToPropagate; ++I)
         if (I->modifiesRegister(DefOp.getReg(), &RI))
           SafeToPropagate = false;
 
-      if (!SafeToPropagate)
+      if (!SafeToPropagate) {
+        DefPoint = Def;
         break;
+      }
 
       DefOp.setIsKill(false);
     }
 
     MachineInstrBuilder Builder =
       BuildMI(MBB, MI, DL, TII.get(AMDGPU::V_ACCVGPR_WRITE_B32_e64), DestReg)
       .add(DefOp);
     if (ImpDefSuperReg)
       Builder.addReg(ImpDefSuperReg, RegState::Define | RegState::Implicit);
 
     if (ImpUseSuperReg) {
       Builder.addReg(ImpUseSuperReg,
                      getKillRegState(KillSrc) | RegState::Implicit);
     }
 
     return;
   }
 
+  // If we were unable to propagate the vgpr register from the defining
+  // accvgpr_write then see if there is an earlier copy that may have inserted
+  // a temporary vgpr register with the same source.
+  for (auto Copy = MI; Copy != DefPoint; ) {
+    --Copy;
+    if (Copy->getOpcode() != AMDGPU::V_ACCVGPR_WRITE_B32_e64)
+      continue;
+
+    MachineOperand &CopyOp = Copy->getOperand(1);
+
+    // Find corresponding v_accvgpr_read that defined the source operand of the
+    // v_accvgpr_write
+    if (CopyOp.isReg()) {
+      for (auto CopyDef = Copy; CopyDef != DefPoint; ) {
+        --CopyDef;
+        if (!CopyDef->definesRegister(CopyOp.getReg(), &RI))
+          continue;
+        if (CopyDef->getOpcode() != AMDGPU::V_ACCVGPR_READ_B32_e64)
+          break;
+
+        // Check if v_accvgpr_read source operand is our agpr source. If so
+        // then this is a vgpr copy that we can use. Check to see if it is safe
+        // to propagate.
+        if (CopyDef->readsRegister(SrcReg)) {
+          bool SafeToPropagate = true;
+          for (auto I = CopyDef; I != MI && SafeToPropagate; ) {
+            ++I;
+            if (I->modifiesRegister(CopyOp.getReg(), &RI))
+              SafeToPropagate = false;
+          }
+
+          if (!SafeToPropagate)
+            break;
+
+          CopyOp.setIsKill(false);
+
+          MachineInstrBuilder Builder =
+            BuildMI(MBB, MI, DL, TII.get(AMDGPU::V_ACCVGPR_WRITE_B32_e64), DestReg)
+            .add(CopyOp);
+          if (ImpDefSuperReg)
+            Builder.addReg(ImpDefSuperReg, RegState::Define | RegState::Implicit);
+
+          if (ImpUseSuperReg) {
+            Builder.addReg(ImpUseSuperReg,
+                           getKillRegState(KillSrc) | RegState::Implicit);
+          }
+
+          return;
+        }
+      }
+    }
+  }
+
   RS.enterBasicBlock(MBB);
   RS.forward(MI);
 
   // Ideally we want to have three registers for a long reg_sequence copy
   // to hide 2 waitstates between v_mov_b32 and accvgpr_write.
   unsigned MaxVGPRs = RI.getRegPressureLimit(&AMDGPU::VGPR_32RegClass,
                                              *MBB.getParent());
 

llvm/test/CodeGen/AMDGPU/accvgpr-copy.mir

--- |
     define amdgpu_kernel void @a16_to_a16() #0 { ret void }
 
     define amdgpu_kernel void @a_to_a_spill() #0 { ret void }
     define amdgpu_kernel void @copy_sgpr_to_agpr_tuple() #0 { ret void }
     define amdgpu_kernel void @copy_sgpr_to_agpr_tuple_kill() #0 { ret void }
     define amdgpu_kernel void @copy_agpr_to_agpr_tuple() #0 { ret void }
     define amdgpu_kernel void @copy_agpr_to_agpr_tuple_kill() #0 { ret void }
 
+    define amdgpu_kernel void @broadcast_a4_with_v0() #0 { ret void }
+    define amdgpu_kernel void @broadcast_a4_unable_to_propagate_v0() #0 { ret void }
+    define amdgpu_kernel void @broadcast_a32_with_v0() #0 { ret void }
+    define amdgpu_kernel void @broadcast_a32_unable_to_propagate_v0() #0 { ret void }
+    define amdgpu_kernel void @broadcast_a32_partially_propagate_v0() #0 { ret void }
+
     attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
 ...
 
 ---
 name:            a_to_v
 tracksRegLiveness: true
 body:             |
   bb.0:
bb.0:
     ; GFX90A: $agpr5 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
     ; GFX90A: $vgpr0 = V_ACCVGPR_READ_B32_e64 killed $agpr0, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3
     ; GFX90A: $agpr4 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec, implicit $exec
     ; GFX90A: S_ENDPGM 0, implicit $agpr4_agpr5_agpr6_agpr7
     S_NOP 0, implicit-def dead $agpr0_agpr1
     renamable $agpr4_agpr5_agpr6_agpr7 = COPY renamable killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
     S_ENDPGM 0, implicit $agpr4_agpr5_agpr6_agpr7
 ...
+
+---
+name: broadcast_a4_with_v0
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $sgpr0_sgpr1
+
+    ; GFX908-LABEL: name: broadcast_a4_with_v0
+    ; GFX908: liveins: $sgpr0_sgpr1
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX908: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX908: $agpr1 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec, implicit $exec
+    ; GFX908: $agpr2 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec, implicit $exec
+    ; GFX908: $agpr3 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec, implicit $exec
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX908: renamable $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX908: $vgpr5 = V_MOV_B32_e32 $sgpr1, implicit $exec, implicit-def $vgpr4_vgpr5, implicit $sgpr0_sgpr1
+    ; GFX908: $vgpr4 = V_MOV_B32_e32 $sgpr0, implicit $exec, implicit killed $sgpr0_sgpr1, implicit $exec
+    ; GFX908: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX908: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX908: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX908: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
+    ; GFX908: GLOBAL_STORE_DWORDX4 killed renamable $vgpr4_vgpr5, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A-LABEL: name: broadcast_a4_with_v0
+    ; GFX90A: liveins: $sgpr0_sgpr1
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX90A: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    ; GFX90A: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX90A: $agpr1 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec, implicit $exec
+    ; GFX90A: $agpr2 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec, implicit $exec
+    ; GFX90A: $agpr3 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec, implicit $exec
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX90A: renamable $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX90A: $vgpr4_vgpr5 = V_PK_MOV_B32 8, $sgpr0_sgpr1, 12, $sgpr0_sgpr1, 0, 0, 0, 0, 0, implicit $exec, implicit killed $sgpr0_sgpr1, implicit $exec
+    ; GFX90A: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX90A: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX90A: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX90A: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
+    ; GFX90A: GLOBAL_STORE_DWORDX4 killed renamable $vgpr4_vgpr5, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $exec :: (store 16, addrspace 1)
+    renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    renamable $agpr1 = COPY renamable $agpr0, implicit $exec
+    renamable $agpr2 = COPY renamable $agpr0, implicit $exec
+    renamable $agpr3 = COPY renamable $agpr0, implicit $exec
+    renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    renamable $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+    renamable $vgpr4_vgpr5 = COPY killed renamable $sgpr0_sgpr1, implicit $exec
+    renamable $vgpr0_vgpr1_vgpr2_vgpr3 = COPY killed renamable $agpr0_agpr1_agpr2_agpr3, implicit $exec
+    GLOBAL_STORE_DWORDX4 killed renamable $vgpr4_vgpr5, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $exec :: (store 16, addrspace 1)
+
+...
+
+---
+name: broadcast_a4_unable_to_propagate_v0
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $sgpr0_sgpr1
+
+    ; GFX908-LABEL: name: broadcast_a4_unable_to_propagate_v0
+    ; GFX908: liveins: $sgpr0_sgpr1
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX908: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX908: renamable $vgpr4 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX908: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX908: $vgpr6 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec
+    ; GFX908: $agpr1 = V_ACCVGPR_WRITE_B32_e64 $vgpr6, implicit $exec, implicit $exec
+    ; GFX908: $agpr2 = V_ACCVGPR_WRITE_B32_e64 $vgpr6, implicit $exec, implicit $exec
+    ; GFX908: $agpr3 = V_ACCVGPR_WRITE_B32_e64 $vgpr6, implicit $exec, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX908: renamable $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr0, killed $vgpr1, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX908: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX908: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX908: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX908: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr4, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: S_ENDPGM 0
+    ; GFX90A-LABEL: name: broadcast_a4_unable_to_propagate_v0
+    ; GFX90A: liveins: $sgpr0_sgpr1
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX90A: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX90A: renamable $vgpr4 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX90A: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX90A: $agpr1 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec, implicit $exec
+    ; GFX90A: $agpr2 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec, implicit $exec
+    ; GFX90A: $agpr3 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX90A: renamable $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr0, killed $vgpr1, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX90A: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX90A: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX90A: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3
+    ; GFX90A: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr4, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: S_ENDPGM 0
+    renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    renamable $vgpr4 = V_MOV_B32_e32 0, implicit $exec
+    renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    renamable $vgpr0 = V_MOV_B32_e32 1065353216, implicit $exec
+    renamable $agpr1 = COPY renamable $agpr0, implicit $exec
+    renamable $agpr2 = COPY renamable $agpr0, implicit $exec
+    renamable $agpr3 = COPY renamable $agpr0, implicit $exec
+    renamable $vgpr1 = V_MOV_B32_e32 1073741824, implicit $exec
+    renamable $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr0, killed $vgpr1, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+    renamable $vgpr0_vgpr1_vgpr2_vgpr3 = COPY killed renamable $agpr0_agpr1_agpr2_agpr3, implicit $exec
+    GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr4, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, addrspace 1)
+    S_ENDPGM 0
+
+...
+
+---
+name: broadcast_a32_with_v0
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $sgpr0_sgpr1
+
+    ; GFX908-LABEL: name: broadcast_a32_with_v0
+    ; GFX908: liveins: $sgpr0_sgpr1
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX908: renamable $sgpr2 = S_MOV_B32 16
+    ; GFX908: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX908: $agpr1 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr2 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr3 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr4 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr5 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr6 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr7 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr8 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr9 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr10 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr11 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr12 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr13 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr14 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr15 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr16 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr17 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr18 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr19 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr20 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr21 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr22 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr23 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr24 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr25 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr26 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr27 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr28 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr29 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr30 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr31 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX908: renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX908: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr4 = V_ACCVGPR_READ_B32_e64 $agpr4, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr5 = V_ACCVGPR_READ_B32_e64 $agpr5, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr6 = V_ACCVGPR_READ_B32_e64 $agpr6, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr7 = V_ACCVGPR_READ_B32_e64 $agpr7, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr8 = V_ACCVGPR_READ_B32_e64 $agpr8, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr9 = V_ACCVGPR_READ_B32_e64 $agpr9, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr10 = V_ACCVGPR_READ_B32_e64 $agpr10, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr11 = V_ACCVGPR_READ_B32_e64 $agpr11, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr12 = V_ACCVGPR_READ_B32_e64 $agpr12, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr13 = V_ACCVGPR_READ_B32_e64 $agpr13, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr14 = V_ACCVGPR_READ_B32_e64 $agpr14, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr15 = V_ACCVGPR_READ_B32_e64 $agpr15, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr16 = V_ACCVGPR_READ_B32_e64 $agpr16, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr17 = V_ACCVGPR_READ_B32_e64 $agpr17, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr18 = V_ACCVGPR_READ_B32_e64 $agpr18, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr19 = V_ACCVGPR_READ_B32_e64 $agpr19, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr20 = V_ACCVGPR_READ_B32_e64 $agpr20, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr21 = V_ACCVGPR_READ_B32_e64 $agpr21, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr22 = V_ACCVGPR_READ_B32_e64 $agpr22, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr23 = V_ACCVGPR_READ_B32_e64 $agpr23, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr24 = V_ACCVGPR_READ_B32_e64 $agpr24, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr25 = V_ACCVGPR_READ_B32_e64 $agpr25, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr26 = V_ACCVGPR_READ_B32_e64 $agpr26, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr27 = V_ACCVGPR_READ_B32_e64 $agpr27, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr28 = V_ACCVGPR_READ_B32_e64 $agpr28, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr29 = V_ACCVGPR_READ_B32_e64 $agpr29, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr30 = V_ACCVGPR_READ_B32_e64 $agpr30, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr31 = V_ACCVGPR_READ_B32_e64 $agpr31, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    ; GFX908: S_ENDPGM 0
+    ; GFX90A-LABEL: name: broadcast_a32_with_v0
+    ; GFX90A: liveins: $sgpr0_sgpr1
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX90A: renamable $sgpr2 = S_MOV_B32 16
+    ; GFX90A: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    ; GFX90A: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX90A: $agpr1 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr2 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr3 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr4 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr5 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr6 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr7 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr8 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr9 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr10 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr11 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr12 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr13 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr14 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr15 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr16 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr17 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr18 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr19 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr20 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr21 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr22 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr23 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr24 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr25 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr26 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr27 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr28 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr29 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr30 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr31 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX90A: renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX90A: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr4 = V_ACCVGPR_READ_B32_e64 $agpr4, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr5 = V_ACCVGPR_READ_B32_e64 $agpr5, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr6 = V_ACCVGPR_READ_B32_e64 $agpr6, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr7 = V_ACCVGPR_READ_B32_e64 $agpr7, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr8 = V_ACCVGPR_READ_B32_e64 $agpr8, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr9 = V_ACCVGPR_READ_B32_e64 $agpr9, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr10 = V_ACCVGPR_READ_B32_e64 $agpr10, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr11 = V_ACCVGPR_READ_B32_e64 $agpr11, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr12 = V_ACCVGPR_READ_B32_e64 $agpr12, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr13 = V_ACCVGPR_READ_B32_e64 $agpr13, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr14 = V_ACCVGPR_READ_B32_e64 $agpr14, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr15 = V_ACCVGPR_READ_B32_e64 $agpr15, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr16 = V_ACCVGPR_READ_B32_e64 $agpr16, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr17 = V_ACCVGPR_READ_B32_e64 $agpr17, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr18 = V_ACCVGPR_READ_B32_e64 $agpr18, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr19 = V_ACCVGPR_READ_B32_e64 $agpr19, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr20 = V_ACCVGPR_READ_B32_e64 $agpr20, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr21 = V_ACCVGPR_READ_B32_e64 $agpr21, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr22 = V_ACCVGPR_READ_B32_e64 $agpr22, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr23 = V_ACCVGPR_READ_B32_e64 $agpr23, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr24 = V_ACCVGPR_READ_B32_e64 $agpr24, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr25 = V_ACCVGPR_READ_B32_e64 $agpr25, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr26 = V_ACCVGPR_READ_B32_e64 $agpr26, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr27 = V_ACCVGPR_READ_B32_e64 $agpr27, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr28 = V_ACCVGPR_READ_B32_e64 $agpr28, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr29 = V_ACCVGPR_READ_B32_e64 $agpr29, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr30 = V_ACCVGPR_READ_B32_e64 $agpr30, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr31 = V_ACCVGPR_READ_B32_e64 $agpr31, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    ; GFX90A: S_ENDPGM 0
+    renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    renamable $sgpr2 = S_MOV_B32 16
+    renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    renamable $agpr1 = COPY renamable $agpr0
+    renamable $agpr2 = COPY renamable $agpr0
+    renamable $agpr3 = COPY renamable $agpr0
+    renamable $agpr4 = COPY renamable $agpr0
+    renamable $agpr5 = COPY renamable $agpr0
+    renamable $agpr6 = COPY renamable $agpr0
+    renamable $agpr7 = COPY renamable $agpr0
+    renamable $agpr8 = COPY renamable $agpr0
+    renamable $agpr9 = COPY renamable $agpr0
+    renamable $agpr10 = COPY renamable $agpr0
+    renamable $agpr11 = COPY renamable $agpr0
+    renamable $agpr12 = COPY renamable $agpr0
+    renamable $agpr13 = COPY renamable $agpr0
+    renamable $agpr14 = COPY renamable $agpr0
+    renamable $agpr15 = COPY renamable $agpr0
+    renamable $agpr16 = COPY renamable $agpr0
+    renamable $agpr17 = COPY renamable $agpr0
+    renamable $agpr18 = COPY renamable $agpr0
+    renamable $agpr19 = COPY renamable $agpr0
+    renamable $agpr20 = COPY renamable $agpr0
+    renamable $agpr21 = COPY renamable $agpr0
+    renamable $agpr22 = COPY renamable $agpr0
+    renamable $agpr23 = COPY renamable $agpr0
+    renamable $agpr24 = COPY renamable $agpr0
+    renamable $agpr25 = COPY renamable $agpr0
+    renamable $agpr26 = COPY renamable $agpr0
+    renamable $agpr27 = COPY renamable $agpr0
+    renamable $agpr28 = COPY renamable $agpr0
+    renamable $agpr29 = COPY renamable $agpr0
+    renamable $agpr30 = COPY renamable $agpr0
+    renamable $agpr31 = COPY renamable $agpr0
+    renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    renamable $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31 = COPY killed renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    S_ENDPGM 0
+
+...
+
+---
+name: broadcast_a32_unable_to_propagate_v0
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $sgpr0_sgpr1
+
+    ; GFX908-LABEL: name: broadcast_a32_unable_to_propagate_v0
+    ; GFX908: liveins: $sgpr0_sgpr1
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX908: renamable $sgpr2 = S_MOV_B32 16
+    ; GFX908: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    ; GFX908: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_MOV_B32_e32 16, implicit $exec
+    ; GFX908: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec
+    ; GFX908: $agpr1 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr2 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr3 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr4 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr5 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr6 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr7 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr8 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr9 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr10 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr11 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr12 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr13 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr14 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr15 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr16 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr17 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr18 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr19 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr20 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr21 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr22 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr23 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr24 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr25 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr26 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr27 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr28 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr29 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr30 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: $agpr31 = V_ACCVGPR_WRITE_B32_e64 $vgpr3, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_ADD_U32_e32 $vgpr0, killed $vgpr1, implicit $exec
+    ; GFX908: GLOBAL_STORE_DWORD_SADDR killed renamable $vgpr0, killed renamable $vgpr1, renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 4, addrspace 1)
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX908: renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX908: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr4 = V_ACCVGPR_READ_B32_e64 $agpr4, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr5 = V_ACCVGPR_READ_B32_e64 $agpr5, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr6 = V_ACCVGPR_READ_B32_e64 $agpr6, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr7 = V_ACCVGPR_READ_B32_e64 $agpr7, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr8 = V_ACCVGPR_READ_B32_e64 $agpr8, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr9 = V_ACCVGPR_READ_B32_e64 $agpr9, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr10 = V_ACCVGPR_READ_B32_e64 $agpr10, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr11 = V_ACCVGPR_READ_B32_e64 $agpr11, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr12 = V_ACCVGPR_READ_B32_e64 $agpr12, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr13 = V_ACCVGPR_READ_B32_e64 $agpr13, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr14 = V_ACCVGPR_READ_B32_e64 $agpr14, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr15 = V_ACCVGPR_READ_B32_e64 $agpr15, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr16 = V_ACCVGPR_READ_B32_e64 $agpr16, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr17 = V_ACCVGPR_READ_B32_e64 $agpr17, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr18 = V_ACCVGPR_READ_B32_e64 $agpr18, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr19 = V_ACCVGPR_READ_B32_e64 $agpr19, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr20 = V_ACCVGPR_READ_B32_e64 $agpr20, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr21 = V_ACCVGPR_READ_B32_e64 $agpr21, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr22 = V_ACCVGPR_READ_B32_e64 $agpr22, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr23 = V_ACCVGPR_READ_B32_e64 $agpr23, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr24 = V_ACCVGPR_READ_B32_e64 $agpr24, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr25 = V_ACCVGPR_READ_B32_e64 $agpr25, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr26 = V_ACCVGPR_READ_B32_e64 $agpr26, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr27 = V_ACCVGPR_READ_B32_e64 $agpr27, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr28 = V_ACCVGPR_READ_B32_e64 $agpr28, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr29 = V_ACCVGPR_READ_B32_e64 $agpr29, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr30 = V_ACCVGPR_READ_B32_e64 $agpr30, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr31 = V_ACCVGPR_READ_B32_e64 $agpr31, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    ; GFX908: S_ENDPGM 0
+    ; GFX90A-LABEL: name: broadcast_a32_unable_to_propagate_v0
+    ; GFX90A: liveins: $sgpr0_sgpr1
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX90A: renamable $sgpr2 = S_MOV_B32 16
+    ; GFX90A: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    ; GFX90A: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_MOV_B32_e32 16, implicit $exec
+    ; GFX90A: $agpr1 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr2 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr3 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr4 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr5 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr6 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr7 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr8 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr9 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr10 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr11 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr12 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr13 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr14 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr15 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr16 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr17 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr18 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr19 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr20 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr21 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr22 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr23 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr24 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr25 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr26 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr27 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr28 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr29 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr30 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr31 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_ADD_U32_e32 $vgpr0, killed $vgpr1, implicit $exec
+    ; GFX90A: GLOBAL_STORE_DWORD_SADDR killed renamable $vgpr0, killed renamable $vgpr1, renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 4, addrspace 1)
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX90A: renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX90A: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr4 = V_ACCVGPR_READ_B32_e64 $agpr4, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr5 = V_ACCVGPR_READ_B32_e64 $agpr5, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr6 = V_ACCVGPR_READ_B32_e64 $agpr6, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr7 = V_ACCVGPR_READ_B32_e64 $agpr7, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr8 = V_ACCVGPR_READ_B32_e64 $agpr8, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr9 = V_ACCVGPR_READ_B32_e64 $agpr9, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr10 = V_ACCVGPR_READ_B32_e64 $agpr10, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr11 = V_ACCVGPR_READ_B32_e64 $agpr11, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr12 = V_ACCVGPR_READ_B32_e64 $agpr12, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr13 = V_ACCVGPR_READ_B32_e64 $agpr13, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr14 = V_ACCVGPR_READ_B32_e64 $agpr14, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr15 = V_ACCVGPR_READ_B32_e64 $agpr15, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr16 = V_ACCVGPR_READ_B32_e64 $agpr16, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr17 = V_ACCVGPR_READ_B32_e64 $agpr17, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr18 = V_ACCVGPR_READ_B32_e64 $agpr18, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr19 = V_ACCVGPR_READ_B32_e64 $agpr19, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr20 = V_ACCVGPR_READ_B32_e64 $agpr20, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr21 = V_ACCVGPR_READ_B32_e64 $agpr21, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr22 = V_ACCVGPR_READ_B32_e64 $agpr22, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr23 = V_ACCVGPR_READ_B32_e64 $agpr23, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr24 = V_ACCVGPR_READ_B32_e64 $agpr24, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr25 = V_ACCVGPR_READ_B32_e64 $agpr25, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr26 = V_ACCVGPR_READ_B32_e64 $agpr26, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr27 = V_ACCVGPR_READ_B32_e64 $agpr27, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr28 = V_ACCVGPR_READ_B32_e64 $agpr28, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr29 = V_ACCVGPR_READ_B32_e64 $agpr29, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr30 = V_ACCVGPR_READ_B32_e64 $agpr30, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr31 = V_ACCVGPR_READ_B32_e64 $agpr31, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    ; GFX90A: S_ENDPGM 0
+    renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    renamable $sgpr2 = S_MOV_B32 16
+    renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    renamable $vgpr0 = V_MOV_B32_e32 0, implicit $exec
+    renamable $vgpr1 = V_MOV_B32_e32 16, implicit $exec
+    renamable $agpr1 = COPY renamable $agpr0
+    renamable $agpr2 = COPY renamable $agpr0
+    renamable $agpr3 = COPY renamable $agpr0
+    renamable $agpr4 = COPY renamable $agpr0
+    renamable $agpr5 = COPY renamable $agpr0
+    renamable $agpr6 = COPY renamable $agpr0
+    renamable $agpr7 = COPY renamable $agpr0
+    renamable $agpr8 = COPY renamable $agpr0
+    renamable $agpr9 = COPY renamable $agpr0
+    renamable $agpr10 = COPY renamable $agpr0
+    renamable $agpr11 = COPY renamable $agpr0
+    renamable $agpr12 = COPY renamable $agpr0
+    renamable $agpr13 = COPY renamable $agpr0
+    renamable $agpr14 = COPY renamable $agpr0
+    renamable $agpr15 = COPY renamable $agpr0
+    renamable $agpr16 = COPY renamable $agpr0
+    renamable $agpr17 = COPY renamable $agpr0
+    renamable $agpr18 = COPY renamable $agpr0
+    renamable $agpr19 = COPY renamable $agpr0
+    renamable $agpr20 = COPY renamable $agpr0
+    renamable $agpr21 = COPY renamable $agpr0
+    renamable $agpr22 = COPY renamable $agpr0
+    renamable $agpr23 = COPY renamable $agpr0
+    renamable $agpr24 = COPY renamable $agpr0
+    renamable $agpr25 = COPY renamable $agpr0
+    renamable $agpr26 = COPY renamable $agpr0
+    renamable $agpr27 = COPY renamable $agpr0
+    renamable $agpr28 = COPY renamable $agpr0
+    renamable $agpr29 = COPY renamable $agpr0
+    renamable $agpr30 = COPY renamable $agpr0
+    renamable $agpr31 = COPY renamable $agpr0
+    renamable $vgpr1 = V_ADD_U32_e32 $vgpr0, killed $vgpr1, implicit $exec
+    GLOBAL_STORE_DWORD_SADDR killed renamable $vgpr0, killed renamable $vgpr1, renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 4, addrspace 1)
+    renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    renamable $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31 = COPY killed renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    S_ENDPGM 0
+
+...
+
+---
+name: broadcast_a32_partially_propagate_v0
+tracksRegLiveness: true
+body:             |
+  bb.0:
+    liveins: $sgpr0_sgpr1
+
+    ; GFX908-LABEL: name: broadcast_a32_partially_propagate_v0
+    ; GFX908: liveins: $sgpr0_sgpr1
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX908: renamable $sgpr2 = S_MOV_B32 16
+    ; GFX908: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX908: $agpr1 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr2 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr3 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr4 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr5 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr6 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr7 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr8 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr9 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: $agpr10 = V_ACCVGPR_WRITE_B32_e64 $vgpr0, implicit $exec
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_MOV_B32_e32 16, implicit $exec
+    ; GFX908: $vgpr4 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec
+    ; GFX908: $agpr11 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr12 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr13 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr14 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr15 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr16 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr17 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr18 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr19 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr20 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr21 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr22 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr23 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr24 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr25 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr26 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr27 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr28 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr29 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr30 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: $agpr31 = V_ACCVGPR_WRITE_B32_e64 $vgpr4, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_ADD_U32_e32 $vgpr0, killed $vgpr1, implicit $exec
+    ; GFX908: GLOBAL_STORE_DWORD_SADDR killed renamable $vgpr0, killed renamable $vgpr1, renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 4, addrspace 1)
+    ; GFX908: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX908: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX908: renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX908: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr4 = V_ACCVGPR_READ_B32_e64 $agpr4, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr5 = V_ACCVGPR_READ_B32_e64 $agpr5, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr6 = V_ACCVGPR_READ_B32_e64 $agpr6, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr7 = V_ACCVGPR_READ_B32_e64 $agpr7, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr8 = V_ACCVGPR_READ_B32_e64 $agpr8, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr9 = V_ACCVGPR_READ_B32_e64 $agpr9, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr10 = V_ACCVGPR_READ_B32_e64 $agpr10, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr11 = V_ACCVGPR_READ_B32_e64 $agpr11, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr12 = V_ACCVGPR_READ_B32_e64 $agpr12, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr13 = V_ACCVGPR_READ_B32_e64 $agpr13, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr14 = V_ACCVGPR_READ_B32_e64 $agpr14, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr15 = V_ACCVGPR_READ_B32_e64 $agpr15, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr16 = V_ACCVGPR_READ_B32_e64 $agpr16, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr17 = V_ACCVGPR_READ_B32_e64 $agpr17, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr18 = V_ACCVGPR_READ_B32_e64 $agpr18, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr19 = V_ACCVGPR_READ_B32_e64 $agpr19, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr20 = V_ACCVGPR_READ_B32_e64 $agpr20, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr21 = V_ACCVGPR_READ_B32_e64 $agpr21, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr22 = V_ACCVGPR_READ_B32_e64 $agpr22, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr23 = V_ACCVGPR_READ_B32_e64 $agpr23, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr24 = V_ACCVGPR_READ_B32_e64 $agpr24, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr25 = V_ACCVGPR_READ_B32_e64 $agpr25, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr26 = V_ACCVGPR_READ_B32_e64 $agpr26, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr27 = V_ACCVGPR_READ_B32_e64 $agpr27, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr28 = V_ACCVGPR_READ_B32_e64 $agpr28, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr29 = V_ACCVGPR_READ_B32_e64 $agpr29, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr30 = V_ACCVGPR_READ_B32_e64 $agpr30, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: $vgpr31 = V_ACCVGPR_READ_B32_e64 $agpr31, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX908: renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX908: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    ; GFX908: S_ENDPGM 0
+    ; GFX90A-LABEL: name: broadcast_a32_partially_propagate_v0
+    ; GFX90A: liveins: $sgpr0_sgpr1
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    ; GFX90A: renamable $sgpr2 = S_MOV_B32 16
+    ; GFX90A: renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    ; GFX90A: renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    ; GFX90A: $agpr1 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr2 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr3 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr4 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr5 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr6 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr7 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr8 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr9 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr10 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_MOV_B32_e32 16, implicit $exec
+    ; GFX90A: $agpr11 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr12 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr13 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr14 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr15 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr16 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr17 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr18 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr19 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr20 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr21 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr22 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr23 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr24 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr25 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr26 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr27 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr28 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr29 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr30 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: $agpr31 = V_ACCVGPR_MOV_B32 $agpr0, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_ADD_U32_e32 $vgpr0, killed $vgpr1, implicit $exec
+    ; GFX90A: GLOBAL_STORE_DWORD_SADDR killed renamable $vgpr0, killed renamable $vgpr1, renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 4, addrspace 1)
+    ; GFX90A: renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    ; GFX90A: renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    ; GFX90A: renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    ; GFX90A: $vgpr0 = V_ACCVGPR_READ_B32_e64 $agpr0, implicit $exec, implicit-def $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr1 = V_ACCVGPR_READ_B32_e64 $agpr1, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr2 = V_ACCVGPR_READ_B32_e64 $agpr2, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr3 = V_ACCVGPR_READ_B32_e64 $agpr3, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr4 = V_ACCVGPR_READ_B32_e64 $agpr4, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr5 = V_ACCVGPR_READ_B32_e64 $agpr5, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr6 = V_ACCVGPR_READ_B32_e64 $agpr6, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr7 = V_ACCVGPR_READ_B32_e64 $agpr7, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr8 = V_ACCVGPR_READ_B32_e64 $agpr8, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr9 = V_ACCVGPR_READ_B32_e64 $agpr9, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr10 = V_ACCVGPR_READ_B32_e64 $agpr10, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr11 = V_ACCVGPR_READ_B32_e64 $agpr11, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr12 = V_ACCVGPR_READ_B32_e64 $agpr12, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr13 = V_ACCVGPR_READ_B32_e64 $agpr13, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr14 = V_ACCVGPR_READ_B32_e64 $agpr14, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr15 = V_ACCVGPR_READ_B32_e64 $agpr15, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr16 = V_ACCVGPR_READ_B32_e64 $agpr16, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr17 = V_ACCVGPR_READ_B32_e64 $agpr17, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr18 = V_ACCVGPR_READ_B32_e64 $agpr18, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr19 = V_ACCVGPR_READ_B32_e64 $agpr19, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr20 = V_ACCVGPR_READ_B32_e64 $agpr20, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr21 = V_ACCVGPR_READ_B32_e64 $agpr21, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr22 = V_ACCVGPR_READ_B32_e64 $agpr22, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr23 = V_ACCVGPR_READ_B32_e64 $agpr23, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr24 = V_ACCVGPR_READ_B32_e64 $agpr24, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr25 = V_ACCVGPR_READ_B32_e64 $agpr25, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr26 = V_ACCVGPR_READ_B32_e64 $agpr26, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr27 = V_ACCVGPR_READ_B32_e64 $agpr27, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr28 = V_ACCVGPR_READ_B32_e64 $agpr28, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr29 = V_ACCVGPR_READ_B32_e64 $agpr29, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr30 = V_ACCVGPR_READ_B32_e64 $agpr30, implicit $exec, implicit $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: $vgpr31 = V_ACCVGPR_READ_B32_e64 $agpr31, implicit $exec, implicit killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    ; GFX90A: renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    ; GFX90A: GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    ; GFX90A: S_ENDPGM 0
+    renamable $vgpr0 = V_MOV_B32_e32 1123418112, implicit $exec
+    renamable $sgpr2 = S_MOV_B32 16
+    renamable $agpr0 = V_ACCVGPR_WRITE_B32_e64 killed $vgpr0, implicit $exec
+    renamable $sgpr0_sgpr1 = S_LOAD_DWORDX2_IMM killed renamable $sgpr0_sgpr1, 36, 0 :: (dereferenceable invariant load 8, align 4, addrspace 4)
+    renamable $agpr1 = COPY renamable $agpr0
+    renamable $agpr2 = COPY renamable $agpr0
+    renamable $agpr3 = COPY renamable $agpr0
+    renamable $agpr4 = COPY renamable $agpr0
+    renamable $agpr5 = COPY renamable $agpr0
+    renamable $agpr6 = COPY renamable $agpr0
+    renamable $agpr7 = COPY renamable $agpr0
+    renamable $agpr8 = COPY renamable $agpr0
+    renamable $agpr9 = COPY renamable $agpr0
+    renamable $agpr10 = COPY renamable $agpr0
+    renamable $vgpr0 = V_MOV_B32_e32 0, implicit $exec
+    renamable $vgpr1 = V_MOV_B32_e32 16, implicit $exec
+    renamable $agpr11 = COPY renamable $agpr0
+    renamable $agpr12 = COPY renamable $agpr0
+    renamable $agpr13 = COPY renamable $agpr0
+    renamable $agpr14 = COPY renamable $agpr0
+    renamable $agpr15 = COPY renamable $agpr0
+    renamable $agpr16 = COPY renamable $agpr0
+    renamable $agpr17 = COPY renamable $agpr0
+    renamable $agpr18 = COPY renamable $agpr0
+    renamable $agpr19 = COPY renamable $agpr0
+    renamable $agpr20 = COPY renamable $agpr0
+    renamable $agpr21 = COPY renamable $agpr0
+    renamable $agpr22 = COPY renamable $agpr0
+    renamable $agpr23 = COPY renamable $agpr0
+    renamable $agpr24 = COPY renamable $agpr0
+    renamable $agpr25 = COPY renamable $agpr0
+    renamable $agpr26 = COPY renamable $agpr0
+    renamable $agpr27 = COPY renamable $agpr0
+    renamable $agpr28 = COPY renamable $agpr0
+    renamable $agpr29 = COPY renamable $agpr0
+    renamable $agpr30 = COPY renamable $agpr0
+    renamable $agpr31 = COPY renamable $agpr0
+    renamable $vgpr1 = V_ADD_U32_e32 $vgpr0, killed $vgpr1, implicit $exec
+    GLOBAL_STORE_DWORD_SADDR killed renamable $vgpr0, killed renamable $vgpr1, renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 4, addrspace 1)
+    renamable $vgpr0 = V_MOV_B32_e32 1073741824, implicit $exec
+    renamable $vgpr1 = V_MOV_B32_e32 1065353216, implicit $exec
+    renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X1F32_e64 $vgpr1, $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
+    renamable $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31 = COPY killed renamable $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31
+    renamable $vgpr32 = V_MOV_B32_e32 0, implicit $exec
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr28_vgpr29_vgpr30_vgpr31, renamable $sgpr0_sgpr1, 112, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr24_vgpr25_vgpr26_vgpr27, renamable $sgpr0_sgpr1, 96, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr20_vgpr21_vgpr22_vgpr23, renamable $sgpr0_sgpr1, 80, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr16_vgpr17_vgpr18_vgpr19, renamable $sgpr0_sgpr1, 64, 0, implicit $exec :: (store 16, align 64, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr12_vgpr13_vgpr14_vgpr15, renamable $sgpr0_sgpr1, 48, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr8_vgpr9_vgpr10_vgpr11, renamable $sgpr0_sgpr1, 32, 0, implicit $exec :: (store 16, align 32, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR renamable $vgpr32, renamable $vgpr4_vgpr5_vgpr6_vgpr7, renamable $sgpr0_sgpr1, 16, 0, implicit $exec :: (store 16, addrspace 1)
+    GLOBAL_STORE_DWORDX4_SADDR killed renamable $vgpr32, killed renamable $vgpr0_vgpr1_vgpr2_vgpr3, killed renamable $sgpr0_sgpr1, 0, 0, implicit $exec :: (store 16, align 128, addrspace 1)
+    S_ENDPGM 0
+
+...

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

bb:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %gep = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %arg, i32 %tid
   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> <float 123.0, float 123.0, float 123.0, float 123.0>, i32 0, i32 0, i32 0)
   ;store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
   store <4 x float> %mai.1, <4 x float> addrspace(1)* %gep
   ret void
 }
 
 ; FIXME: Resulting code for splat is pretty bad. A v_mov_b32 is moved
 ; in the middle of the expanded agpr reg_sequence. The broadcast of
 ; the individual AGPR->AGPR components should avoid the intermediate AGPR case.
 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_lit_splat_bad_code:
 ; GFX908_A: v_mov_b32_e32 [[TMP0:v[0-9]+]], 0x42f60000
 ; GCN:      v_accvgpr_write_b32 [[AGPR:a[0-9]+]], [[TMP0]]
 ; GFX908:   s_nop 0
 ; GFX908:   v_accvgpr_read_b32 [[TMP1:v[0-9]+]], [[AGPR]]
-; GFX908:   v_accvgpr_read_b32 [[TMP2:v[0-9]+]], [[AGPR]]
-; GFX908:   v_accvgpr_read_b32 [[TMP3:v[0-9]+]], [[AGPR]]
+; GFX908:   s_nop 1
+; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
 ; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
 ; GFX90A-COUNT-3: v_accvgpr_mov_b32 a{{[0-9]+}}, [[AGPR]]
 ; GCN: s_nop 0
 ; GFX908_A:  v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
 ; GFX908-COUNT-4: v_accvgpr_read_b32
 ; GFX908:    global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}]
 ; GFX90A:    global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}], s[{{[0-9:]+}}]
 define amdgpu_kernel void @test_mfma_f32_4x4x1f32_lit_splat_bad_code(<4 x float> addrspace(1)* %arg) {
 bb: