diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp --- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp @@ -508,6 +508,7 @@ 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; @@ -527,8 +528,10 @@ if (I->modifiesRegister(DefOp.getReg(), &RI)) SafeToPropagate = false; - if (!SafeToPropagate) + if (!SafeToPropagate) { + DefPoint = Def; break; + } DefOp.setIsKill(false); } @@ -547,6 +550,59 @@ 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); diff --git a/llvm/test/CodeGen/AMDGPU/accvgpr-copy.mir b/llvm/test/CodeGen/AMDGPU/accvgpr-copy.mir --- a/llvm/test/CodeGen/AMDGPU/accvgpr-copy.mir +++ b/llvm/test/CodeGen/AMDGPU/accvgpr-copy.mir @@ -41,6 +41,12 @@ 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" } ... @@ -1106,3 +1112,815 @@ 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 + +... diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll @@ -586,11 +586,10 @@ ; 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]+}}]