Skip to content

AMDGPU: Remove unnecessary AGPR operand legalization #162093

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Oct 7, 2025
+930 −1,141
Merged

AMDGPU: Remove unnecessary AGPR operand legalization #162093

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
25ad7cd
AMDGPU: Fix mfma agpr allocation failures with -O0
arsenm Oct 5, 2025
764d91d
AMDGPU: Track minNumAGPRs in MFI instead of mayUseAGPRs
arsenm Oct 5, 2025
c94ce69
AMDGPU: Remove unnecessary AGPR operand legalization
arsenm Oct 5, 2025
758b376
Merge branch 'main' into users/arsenm/amdgpu/remove-unnecessary-agpr-…
arsenm Oct 6, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
77 changes: 13 additions & 64 deletions llvm/lib/Target/AMDGPU/SIISelLowering.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17346,75 +17346,24 @@ void SITargetLowering::AdjustInstrPostInstrSelection(MachineInstr &MI,

MachineFunction *MF = MI.getParent()->getParent();
MachineRegisterInfo &MRI = MF->getRegInfo();
SIMachineFunctionInfo *Info = MF->getInfo<SIMachineFunctionInfo>();

if (TII->isVOP3(MI.getOpcode())) {
// Make sure constant bus requirements are respected.
TII->legalizeOperandsVOP3(MRI, MI);

// Prefer VGPRs over AGPRs in mAI instructions where possible.
// This saves a chain-copy of registers and better balance register
// use between vgpr and agpr as agpr tuples tend to be big.
if (!MI.getDesc().operands().empty()) {
unsigned Opc = MI.getOpcode();
bool HasAGPRs =
!Subtarget->hasGFX90AInsts() || Info->getMinNumAGPRs() != 0;
const SIRegisterInfo *TRI = Subtarget->getRegisterInfo();
int16_t Src2Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src2);
for (auto I :
{AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0),
AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src1), Src2Idx}) {
if (I == -1)
break;
if ((I == Src2Idx) && (HasAGPRs))
break;
MachineOperand &Op = MI.getOperand(I);
if (!Op.isReg() || !Op.getReg().isVirtual())
continue;
auto *RC = TRI->getRegClassForReg(MRI, Op.getReg());
if (!TRI->hasAGPRs(RC))
continue;
auto *Src = MRI.getUniqueVRegDef(Op.getReg());
if (!Src || !Src->isCopy() ||
!TRI->isSGPRReg(MRI, Src->getOperand(1).getReg()))
continue;
auto *NewRC = TRI->getEquivalentVGPRClass(RC);
// All uses of agpr64 and agpr32 can also accept vgpr except for
// v_accvgpr_read, but we do not produce agpr reads during selection,
// so no use checks are needed.
MRI.setRegClass(Op.getReg(), NewRC);
}

if (TII->isMAI(MI)) {
// The ordinary src0, src1, src2 were legalized above.
//
// We have to also legalize the appended v_mfma_ld_scale_b32 operands,
// as a separate instruction.
int Src0Idx = AMDGPU::getNamedOperandIdx(MI.getOpcode(),
AMDGPU::OpName::scale_src0);
if (Src0Idx != -1) {
int Src1Idx = AMDGPU::getNamedOperandIdx(MI.getOpcode(),
AMDGPU::OpName::scale_src1);
if (TII->usesConstantBus(MRI, MI, Src0Idx) &&
TII->usesConstantBus(MRI, MI, Src1Idx))
TII->legalizeOpWithMove(MI, Src1Idx);
}
}

if (!HasAGPRs)
return;

// Resolve the rest of AV operands to AGPRs.
if (auto *Src2 = TII->getNamedOperand(MI, AMDGPU::OpName::src2)) {
if (Src2->isReg() && Src2->getReg().isVirtual()) {
auto *RC = TRI->getRegClassForReg(MRI, Src2->getReg());
if (TRI->isVectorSuperClass(RC)) {
auto *NewRC = TRI->getEquivalentAGPRClass(RC);
MRI.setRegClass(Src2->getReg(), NewRC);
if (Src2->isTied())
MRI.setRegClass(MI.getOperand(0).getReg(), NewRC);
}
}
if (TII->isMAI(MI)) {
// The ordinary src0, src1, src2 were legalized above.
//
// We have to also legalize the appended v_mfma_ld_scale_b32 operands,
// as a separate instruction.
int Src0Idx = AMDGPU::getNamedOperandIdx(MI.getOpcode(),
AMDGPU::OpName::scale_src0);
if (Src0Idx != -1) {
int Src1Idx = AMDGPU::getNamedOperandIdx(MI.getOpcode(),
AMDGPU::OpName::scale_src1);
if (TII->usesConstantBus(MRI, MI, Src0Idx) &&
TII->usesConstantBus(MRI, MI, Src1Idx))
TII->legalizeOpWithMove(MI, Src1Idx);
}
}

Expand Down
80 changes: 40 additions & 40 deletions llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx90a.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -726,12 +726,12 @@ define amdgpu_kernel void @test_mfma_f64_4x4x4f64(ptr addrspace(1) %arg, double
; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: s_nop 1
; GFX90A-VGPR-NEXT: v_mfma_f64_4x4x4f64 v[4:5], v[0:1], v[2:3], 0
; GFX90A-VGPR-NEXT: v_mfma_f64_4x4x4f64 v[0:1], v[2:3], v[4:5], 0
; GFX90A-VGPR-NEXT: s_nop 3
; GFX90A-VGPR-NEXT: v_mfma_f64_4x4x4f64 v[0:1], v[0:1], v[2:3], v[4:5] cbsz:1 abid:2 blgp:3
; GFX90A-VGPR-NEXT: v_mfma_f64_4x4x4f64 v[0:1], v[2:3], v[4:5], v[0:1] cbsz:1 abid:2 blgp:3
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, 0
; GFX90A-VGPR-NEXT: s_nop 7
; GFX90A-VGPR-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1]
Expand All @@ -742,12 +742,12 @@ define amdgpu_kernel void @test_mfma_f64_4x4x4f64(ptr addrspace(1) %arg, double
; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[2:3]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[6:7]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[6:7]
; GFX942-VGPR-NEXT: s_nop 1
; GFX942-VGPR-NEXT: v_mfma_f64_4x4x4_4b_f64 v[4:5], v[0:1], v[2:3], 0
; GFX942-VGPR-NEXT: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[2:3], v[4:5], 0
; GFX942-VGPR-NEXT: s_nop 3
; GFX942-VGPR-NEXT: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], v[2:3], v[4:5] cbsz:1 abid:2 neg:[1,1,0]
; GFX942-VGPR-NEXT: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[2:3], v[4:5], v[0:1] cbsz:1 abid:2 neg:[1,1,0]
; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, 0
; GFX942-VGPR-NEXT: s_nop 7
; GFX942-VGPR-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1]
Expand All @@ -765,10 +765,10 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
; GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24
; GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-NEXT: v_mov_b32_e32 v2, s10
; GFX90A-NEXT: v_mov_b32_e32 v0, s10
; GFX90A-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0
; GFX90A-NEXT: v_mov_b32_e32 v3, s11
; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[12:13], s[12:13] op_sel:[0,1]
; GFX90A-NEXT: v_mov_b32_e32 v1, s11
; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[12:13], s[12:13] op_sel:[0,1]
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0
; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1
Expand All @@ -779,7 +779,7 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6
; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7
; GFX90A-NEXT: s_nop 1
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[2:3], v[0:1], a[0:7] cbsz:1 abid:2 blgp:3
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:3
; GFX90A-NEXT: v_mov_b32_e32 v0, 0
; GFX90A-NEXT: s_nop 15
; GFX90A-NEXT: s_nop 0
Expand All @@ -792,10 +792,10 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
; GFX942-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24
; GFX942-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-NEXT: v_mov_b32_e32 v2, s10
; GFX942-NEXT: v_mov_b32_e32 v0, s10
; GFX942-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0
; GFX942-NEXT: v_mov_b32_e32 v3, s11
; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
; GFX942-NEXT: v_mov_b32_e32 v1, s11
; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[12:13]
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-NEXT: v_accvgpr_write_b32 a0, s0
; GFX942-NEXT: v_accvgpr_write_b32 a1, s1
Expand All @@ -806,7 +806,7 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
; GFX942-NEXT: v_accvgpr_write_b32 a6, s6
; GFX942-NEXT: v_accvgpr_write_b32 a7, s7
; GFX942-NEXT: s_nop 1
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[2:3], v[0:1], a[0:7] cbsz:1 abid:2 neg:[1,1,0]
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]
; GFX942-NEXT: v_mov_b32_e32 v0, 0
; GFX942-NEXT: s_nop 15
; GFX942-NEXT: s_nop 0
Expand All @@ -819,17 +819,17 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
; GFX90A-VGPR-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24
; GFX90A-VGPR-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34
; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v10, s10
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, s10
; GFX90A-VGPR-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v11, s11
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[12:13], s[12:13] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v9, s11
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[12:13], s[12:13] op_sel:[0,1]
; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], s[4:5], s[4:5] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: s_nop 1
; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[10:11], v[8:9], v[0:7] cbsz:1 abid:2 blgp:3
; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:3
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 0
; GFX90A-VGPR-NEXT: s_nop 15
; GFX90A-VGPR-NEXT: s_nop 0
Expand All @@ -842,17 +842,17 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
; GFX942-VGPR-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24
; GFX942-VGPR-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34
; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-VGPR-NEXT: v_mov_b32_e32 v10, s10
; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, s10
; GFX942-VGPR-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0
; GFX942-VGPR-NEXT: v_mov_b32_e32 v11, s11
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[12:13]
; GFX942-VGPR-NEXT: v_mov_b32_e32 v9, s11
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[12:13]
; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[4:5]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[6:7]
; GFX942-VGPR-NEXT: s_nop 1
; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[10:11], v[8:9], v[0:7] cbsz:1 abid:2 neg:[1,1,0]
; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]
; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 0
; GFX942-VGPR-NEXT: s_nop 15
; GFX942-VGPR-NEXT: s_nop 0
Expand Down Expand Up @@ -1629,20 +1629,20 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_imm(ptr addrspace(1) %arg, d
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, 0x3ff00000
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, v0
; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v12, s2
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v13, s3
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v10, s2
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v11, s3
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, v0
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, v0
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, v0
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, v0
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, v0
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], v[6:7], v[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[12:13], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], v[4:5], v[4:5] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], v[2:3], v[2:3] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], v[0:1], v[0:1] op_sel:[0,1]
; GFX90A-VGPR-NEXT: s_nop 1
; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[12:13], v[10:11], v[2:9]
; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[10:11], v[12:13], v[2:9]
; GFX90A-VGPR-NEXT: s_nop 15
; GFX90A-VGPR-NEXT: s_nop 1
; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:16
Expand All @@ -1657,20 +1657,20 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_imm(ptr addrspace(1) %arg, d
; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, 0x3ff00000
; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, v0
; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-VGPR-NEXT: v_mov_b32_e32 v12, s2
; GFX942-VGPR-NEXT: v_mov_b32_e32 v13, s3
; GFX942-VGPR-NEXT: v_mov_b32_e32 v10, s2
; GFX942-VGPR-NEXT: v_mov_b32_e32 v11, s3
; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, v0
; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, v0
; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, v0
; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, v0
; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, v0
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], v[6:7]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[6:7]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], v[4:5]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], v[2:3]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], v[0:1]
; GFX942-VGPR-NEXT: s_nop 1
; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[12:13], v[10:11], v[2:9]
; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[10:11], v[12:13], v[2:9]
; GFX942-VGPR-NEXT: s_nop 15
; GFX942-VGPR-NEXT: s_nop 1
; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:16
Expand Down Expand Up @@ -1743,20 +1743,20 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_lit(ptr addrspace(1) %
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, 0x405ec000
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, v0
; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v12, s2
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v13, s3
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v10, s2
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v11, s3
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, v1
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, v0
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, v1
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, v0
; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, v1
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], v[6:7], v[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[12:13], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], v[4:5], v[4:5] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], v[2:3], v[2:3] op_sel:[0,1]
; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], v[0:1], v[0:1] op_sel:[0,1]
; GFX90A-VGPR-NEXT: s_nop 1
; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[12:13], v[10:11], v[2:9]
; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[10:11], v[12:13], v[2:9]
; GFX90A-VGPR-NEXT: s_nop 15
; GFX90A-VGPR-NEXT: s_nop 1
; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:16
Expand All @@ -1771,20 +1771,20 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_lit(ptr addrspace(1) %
; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, 0x405ec000
; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, v0
; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-VGPR-NEXT: v_mov_b32_e32 v12, s2
; GFX942-VGPR-NEXT: v_mov_b32_e32 v13, s3
; GFX942-VGPR-NEXT: v_mov_b32_e32 v10, s2
; GFX942-VGPR-NEXT: v_mov_b32_e32 v11, s3
; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, v1
; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, v0
; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, v1
; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, v0
; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, v1
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], v[6:7]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[6:7]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], v[4:5]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], v[2:3]
; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], v[0:1]
; GFX942-VGPR-NEXT: s_nop 1
; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[12:13], v[10:11], v[2:9]
; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[10:11], v[12:13], v[2:9]
; GFX942-VGPR-NEXT: s_nop 15
; GFX942-VGPR-NEXT: s_nop 1
; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:16
Expand Down
Loading