[clang] [llvm] AMDGPU: Define v_mfma_f32_{16x16x128|32x32x64}_f8f6f4 instructions (PR #116723)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Thu Nov 21 12:52:53 PST 2024
================
@@ -15454,6 +15454,23 @@ void SITargetLowering::AdjustInstrPostInstrSelection(MachineInstr &MI,
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 = Src0Idx + 2;
+ assert(Src1Idx = AMDGPU::getNamedOperandIdx(
----------------
arsenm wrote:
yes
https://github.com/llvm/llvm-project/pull/116723
More information about the cfe-commits
mailing list