[PATCH] D122337: [AMDGPU] Use GenericTable to classify DGEMM

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

rampitec created this revision.
rampitec added reviewers: Joe_Nash, kzhuravl.
Herald added subscribers: foad, kerbowa, hiraditya, t-tye, tpr, dstuttard, yaxunl, nhaehnle, jvesely, arsenm.
Herald added a project: All.
rampitec requested review of this revision.
Herald added a subscriber: wdng.
Herald added a project: LLVM.

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


https://reviews.llvm.org/D122337

Files:
  llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
  llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
  llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
  llvm/lib/Target/AMDGPU/VOP3PInstructions.td


Index: llvm/lib/Target/AMDGPU/VOP3PInstructions.td
===================================================================
--- llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -487,6 +487,7 @@
 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;
 }
 
@@ -559,8 +560,10 @@
   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 {
@@ -590,7 +593,7 @@
   let FilterClass = "MAIInst";
   let CppTypeName = "MAIInstInfo";
   let Fields = [
-    "Opcode", "is_gfx940_xdl"
+    "Opcode", "is_dgemm", "is_gfx940_xdl"
   ];
 
   let PrimaryKey = ["Opcode"];
Index: llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
===================================================================
--- llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -70,6 +70,7 @@
 
 struct MAIInstInfo {
   uint16_t Opcode;
+  bool is_dgemm;
   bool is_gfx940_xdl;
 };
 
@@ -450,6 +451,9 @@
 LLVM_READONLY
 bool getVOP3IsSingle(unsigned Opc);
 
+LLVM_READONLY
+bool getMAIIsDGEMM(unsigned Opc);
+
 LLVM_READONLY
 bool getMAIIsGFX940XDL(unsigned Opc);
 
Index: llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -343,6 +343,11 @@
   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;
Index: llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -92,12 +92,7 @@
 }
 
 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) {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D122337.417699.patch
Type: text/x-patch
Size: 3052 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20220323/c908bbd6/attachment.bin>