[AMDGPU] Use GenericTable to classify DGEMM

[rampitec]

Summary

Since there is a table introduced for MAI instructions extend it
to use for DGEMM classification.

[foad]

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h

455

Can you add a brief comment here (or somewhere) saying roughly what DGEMM means here? Especially since there is no public documentation.

[rampitec]

Added comment.

llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp

static bool isSMovRel(unsigned Opcode) {
   case AMDGPU::S_MOVRELD_B64:
     return true;
   default:
     return false;
   }
 }
 
 static bool isDGEMM(unsigned Opcode) {
-  return Opcode == AMDGPU::V_MFMA_F64_4X4X4F64_e64 ||
-         Opcode == AMDGPU::V_MFMA_F64_4X4X4F64_vgprcd_e64 ||
-         Opcode == AMDGPU::V_MFMA_F64_16X16X4F64_e64 ||
-         Opcode == AMDGPU::V_MFMA_F64_16X16X4F64_vgprcd_e64 ||
-         Opcode == AMDGPU::V_MFMA_F64_16X16X4F64_mac_e64 ||
-         Opcode == AMDGPU::V_MFMA_F64_16X16X4F64_mac_vgprcd_e64;
+  return AMDGPU::getMAIIsDGEMM(Opcode);
 }
 
 static bool isXDL(const GCNSubtarget &ST, const MachineInstr &MI) {
   unsigned Opcode = MI.getOpcode();
 
   if (!SIInstrInfo::isMAI(MI) ||
       isDGEMM(Opcode) ||
       Opcode == AMDGPU::V_ACCVGPR_WRITE_B32_e64 ||
       Opcode == AMDGPU::V_ACCVGPR_READ_B32_e64)
     return false;
 
   if (!ST.hasGFX940Insts())
     return true;
 
   return AMDGPU::getMAIIsGFX940XDL(Opcode);
 }
 
 static bool isSendMsgTraceDataOrGDS(const SIInstrInfo &TII,
                                     const MachineInstr &MI) {
   if (TII.isAlwaysGDS(MI.getOpcode()))
     return true;
 
   switch (MI.getOpcode()) {
   case AMDGPU::S_SENDMSG:
   case AMDGPU::S_SENDMSGHALT:
   case AMDGPU::S_TTRACEDATA:
     return true;
   // These DS opcodes don't support GDS.
   case AMDGPU::DS_NOP:
   case AMDGPU::DS_PERMUTE_B32:
   case AMDGPU::DS_BPERMUTE_B32:

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h

struct GcnBufferFormatInfo {
   unsigned BitsPerComp;
   unsigned NumComponents;
   unsigned NumFormat;
   unsigned DataFormat;
 };
 
 struct MAIInstInfo {
   uint16_t Opcode;
+  bool is_dgemm;
   bool is_gfx940_xdl;
 };
 
 #define GET_MIMGBaseOpcode_DECL
 #define GET_MIMGDim_DECL
 #define GET_MIMGEncoding_DECL
 #define GET_MIMGLZMapping_DECL
 #define GET_MIMGMIPMapping_DECL

 bool getVOP1IsSingle(unsigned Opc);
 
 LLVM_READONLY
 bool getVOP2IsSingle(unsigned Opc);
 
 LLVM_READONLY
 bool getVOP3IsSingle(unsigned Opc);
 
+/// Returns true if MAI operation is a double precision GEMM.
+LLVM_READONLY
+bool getMAIIsDGEMM(unsigned Opc);
+
 LLVM_READONLY
 bool getMAIIsGFX940XDL(unsigned Opc);
 
 LLVM_READONLY
 const GcnBufferFormatInfo *getGcnBufferFormatInfo(uint8_t BitsPerComp,
                                                   uint8_t NumComponents,
                                                   uint8_t NumFormat,
                                                   const MCSubtargetInfo &STI);

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

bool getVOP2IsSingle(unsigned Opc) {
   return Info ? Info->IsSingle : false;
 }
 
 bool getVOP3IsSingle(unsigned Opc) {
   const VOPInfo *Info = getVOP3OpcodeHelper(Opc);
   return Info ? Info->IsSingle : false;
 }
 
+bool getMAIIsDGEMM(unsigned Opc) {
+  const MAIInstInfo *Info = getMAIInstInfoHelper(Opc);
+  return Info ? Info->is_dgemm : false;
+}
+
 bool getMAIIsGFX940XDL(unsigned Opc) {
   const MAIInstInfo *Info = getMAIInstInfoHelper(Opc);
   return Info ? Info->is_gfx940_xdl : false;
 }
 
 // Wrapper for Tablegen'd function.  enum Subtarget is not defined in any
 // header files, so we need to wrap it in a function that takes unsigned
 // instead.

llvm/lib/Target/AMDGPU/VOP3PInstructions.td

let isAsCheapAsAMove = 1, isReMaterializable = 1 in {
   let isMoveImm = 1 in {
     defm V_ACCVGPR_WRITE_B32 : VOP3Inst<"v_accvgpr_write_b32", VOPProfileAccWrite>;
   } // End isMoveImm = 1
 } // End isAsCheapAsAMove = 1, isReMaterializable = 1
 
 class MAIInst<string OpName, VOPProfile P, SDPatternOperator node>
   : VOP3InstBase<OpName, P, node> {
   Instruction Opcode = !cast<Instruction>(NAME);
+  bit is_dgemm = 0;
   bit is_gfx940_xdl = 0;
 }
 
 multiclass MAIInst<string OpName, string P, SDPatternOperator node,
                    bit NoDstOverlap = !cast<VOPProfileMAI>("VOPProfileMAI_" # P).NoDstOverlap> {
   let isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 in {
     // FP32 denorm mode is respected, rounding mode is not. Exceptions are not supported.
     let Constraints = !if(NoDstOverlap, "@earlyclobber $vdst", "") in {
let Predicates = [isGFX90APlus] in {
   let is_gfx940_xdl = 1 in {
   defm V_MFMA_F32_32X32X4BF16_1K  : MAIInst<"v_mfma_f32_32x32x4bf16_1k",  "F32_V4I16_X32",  int_amdgcn_mfma_f32_32x32x4bf16_1k>;
   defm V_MFMA_F32_16X16X4BF16_1K  : MAIInst<"v_mfma_f32_16x16x4bf16_1k",  "F32_V4I16_X16",  int_amdgcn_mfma_f32_16x16x4bf16_1k>;
   defm V_MFMA_F32_4X4X4BF16_1K    : MAIInst<"v_mfma_f32_4x4x4bf16_1k",    "F32_V4I16_X4",   int_amdgcn_mfma_f32_4x4x4bf16_1k>;
   defm V_MFMA_F32_32X32X8BF16_1K  : MAIInst<"v_mfma_f32_32x32x8bf16_1k",  "F32_V4I16_X16",  int_amdgcn_mfma_f32_32x32x8bf16_1k>;
   defm V_MFMA_F32_16X16X16BF16_1K : MAIInst<"v_mfma_f32_16x16x16bf16_1k", "F32_V4I16_X4",   int_amdgcn_mfma_f32_16x16x16bf16_1k>;
   }
 
+  let is_dgemm = 1 in {
   defm V_MFMA_F64_16X16X4F64      : MAIInst<"v_mfma_f64_16x16x4f64",      "F64_16X16X4F64", int_amdgcn_mfma_f64_16x16x4f64>;
   defm V_MFMA_F64_4X4X4F64        : MAIInst<"v_mfma_f64_4x4x4f64",        "F64_4X4X4F64",   int_amdgcn_mfma_f64_4x4x4f64>;
+  }
 } // End Predicates = [isGFX90APlus]
 
 let Predicates = [isGFX940Plus], is_gfx940_xdl = 1 in {
   defm V_MFMA_I32_32X32X16I8       : MAIInst<"v_mfma_i32_32x32x16i8",       "I32_I64_X32",    int_amdgcn_mfma_i32_32x32x16_i8>;
   defm V_MFMA_I32_16X16X32I8       : MAIInst<"v_mfma_i32_16x16x32i8",       "I32_I64_X16",    int_amdgcn_mfma_i32_16x16x32_i8>;
   defm V_MFMA_F32_16X16X8XF32      : MAIInst<"v_mfma_f32_16x16x8xf32",      "F32_V2F32_X16",  int_amdgcn_mfma_f32_16x16x8_xf32>;
   defm V_MFMA_F32_32X32X4XF32      : MAIInst<"v_mfma_f32_32x32x4xf32",      "F32_V2F32_X32",  int_amdgcn_mfma_f32_32x32x4_xf32>;
 } // End Predicates = [isGFX940Plus], is_gfx940_xdl = 1

 defm V_SMFMAC_I32_16X16X64_I8      : SMFMACInst<"v_smfmac_i32_16x16x64_i8",      "I32_16X16X64_I8",  int_amdgcn_smfmac_i32_16x16x64_i8>;
 defm V_SMFMAC_I32_32X32X32_I8      : SMFMACInst<"v_smfmac_i32_32x32x32_i8",      "I32_32X32X32_I8",  int_amdgcn_smfmac_i32_32x32x32_i8>;
 }
 
 def MAIInstInfoTable : GenericTable {
   let FilterClass = "MAIInst";
   let CppTypeName = "MAIInstInfo";
   let Fields = [
-    "Opcode", "is_gfx940_xdl"
+    "Opcode", "is_dgemm", "is_gfx940_xdl"
   ];
 
   let PrimaryKey = ["Opcode"];
   let PrimaryKeyName = "getMAIInstInfoHelper";
 }
 
 let SubtargetPredicate = HasPackedFP32Ops, isCommutable = 1 in {
   defm V_PK_FMA_F32 : VOP3PInst<"v_pk_fma_f32", VOP3_Profile<VOP_V2F32_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fma>;