[clang] [llvm] AMDGPU: Define v_mfma_f32_{16x16x128|32x32x64}_f8f6f4 instructions (PR #116723)

Mariusz Sikora via cfe-commits cfe-commits at lists.llvm.org
Thu Nov 21 12:51:54 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(
----------------
mariusz-sikora-at-amd wrote:

== ?

https://github.com/llvm/llvm-project/pull/116723


More information about the cfe-commits mailing list