[llvm] [AMDGPU] Replace tablegen field is_gfx940_xdl with is_sgemm. NFC. (PR #143865)
Jay Foad via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 12 02:51:39 PDT 2025
https://github.com/jayfoad created https://github.com/llvm/llvm-project/pull/143865
This is an attempt to make the tablegen files declare instruction
properties that make sense regardless of subtarget. The GFX940-specific
behavior is explained and implemented in SIInstrInfo::isXDL.
>From 948f22726890ba9915e5c8bcd7aa1b1abb183b88 Mon Sep 17 00:00:00 2001
From: Jay Foad <jay.foad at amd.com>
Date: Thu, 12 Jun 2025 10:46:45 +0100
Subject: [PATCH] [AMDGPU] Replace tablegen field is_gfx940_xdl with is_sgemm.
NFC.
This is an attempt to make the tablegen files declare instruction
properties that make sense regardless of subtarget. The GFX940-specific
behavior is explained and implemented in SIInstrInfo::isXDL.
---
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 7 +++--
llvm/lib/Target/AMDGPU/SIInstrInfo.h | 2 ++
.../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 4 +--
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 5 ++--
llvm/lib/Target/AMDGPU/VOP3PInstructions.td | 28 +++++++++----------
5 files changed, 25 insertions(+), 21 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index 2ebf8b99e9d7b..963ea2ffe97b2 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -10371,8 +10371,9 @@ bool SIInstrInfo::isXDL(const MachineInstr &MI) const {
Opcode == AMDGPU::V_ACCVGPR_READ_B32_e64)
return false;
- if (!ST.hasGFX940Insts())
- return true;
+ // On GFX940+ XDL does not include SGEMM instructions.
+ if (ST.hasGFX940Insts() && isSGEMM(Opcode))
+ return false;
- return AMDGPU::getMAIIsGFX940XDL(Opcode);
+ return true;
}
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 01dd3c9f4119e..9b80b3b486e4f 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -867,6 +867,8 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
static bool isDGEMM(unsigned Opcode) { return AMDGPU::getMAIIsDGEMM(Opcode); }
+ static bool isSGEMM(unsigned Opcode) { return AMDGPU::getMAIIsSGEMM(Opcode); }
+
static bool isLDSDIR(const MachineInstr &MI) {
return MI.getDesc().TSFlags & SIInstrFlags::LDSDIR;
}
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index c0cd43a9c35df..817e8facd251e 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -556,9 +556,9 @@ bool getMAIIsDGEMM(unsigned Opc) {
return Info && Info->is_dgemm;
}
-bool getMAIIsGFX940XDL(unsigned Opc) {
+bool getMAIIsSGEMM(unsigned Opc) {
const MAIInstInfo *Info = getMAIInstInfoHelper(Opc);
- return Info && Info->is_gfx940_xdl;
+ return Info && Info->is_sgemm;
}
uint8_t mfmaScaleF8F6F4FormatToNumRegs(unsigned EncodingVal) {
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
index 975a8908059c1..991bd540d4e3c 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -99,7 +99,7 @@ struct GcnBufferFormatInfo {
struct MAIInstInfo {
uint16_t Opcode;
bool is_dgemm;
- bool is_gfx940_xdl;
+ bool is_sgemm;
};
struct MFMA_F8F6F4_Info {
@@ -585,8 +585,9 @@ bool isVOPCAsmOnly(unsigned Opc);
LLVM_READONLY
bool getMAIIsDGEMM(unsigned Opc);
+/// Returns true if MAI operation is a single precision (non xf32) GEMM.
LLVM_READONLY
-bool getMAIIsGFX940XDL(unsigned Opc);
+bool getMAIIsSGEMM(unsigned Opc);
struct CanBeVOPD {
bool X;
diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
index e8db879ca5077..1f7c2a4e83101 100644
--- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -880,7 +880,7 @@ class MAIInst<string OpName, VOPProfile P, SDPatternOperator node, bit Scaled =
let SubtargetPredicate = HasMAIInsts;
Instruction Opcode = !cast<Instruction>(NAME);
bit is_dgemm = 0;
- bit is_gfx940_xdl = 0;
+ bit is_sgemm = 1;
let PseudoInstr = NAME; // FIXME: Why is this not the default
}
@@ -1005,7 +1005,7 @@ defm V_MFMA_F32_16X16X4F32 : MAIInst<"v_mfma_f32_16x16x4f32", "F32_F32_X4",
defm V_MFMA_F32_32X32X1F32 : MAIInst<"v_mfma_f32_32x32x1f32", "F32_F32_X32", int_amdgcn_mfma_f32_32x32x1f32>;
defm V_MFMA_F32_32X32X2F32 : MAIInst<"v_mfma_f32_32x32x2f32", "F32_F32_X16", int_amdgcn_mfma_f32_32x32x2f32>;
-let is_gfx940_xdl = 1 in {
+let is_sgemm = 0 in {
defm V_MFMA_F32_4X4X4F16 : MAIInst<"v_mfma_f32_4x4x4f16", "F32_V4F16_X4", int_amdgcn_mfma_f32_4x4x4f16>;
defm V_MFMA_I32_4X4X4I8 : MAIInst<"v_mfma_i32_4x4x4i8", "I32_I32_X4", int_amdgcn_mfma_i32_4x4x4i8>;
defm V_MFMA_F32_16X16X4F16 : MAIInst<"v_mfma_f32_16x16x4f16", "F32_V4F16_X16", int_amdgcn_mfma_f32_16x16x4f16>;
@@ -1026,7 +1026,7 @@ defm V_MFMA_F32_32X32X2BF16 : MAIInst<"v_mfma_f32_32x32x2bf16", "F32_V2I16_X32",
defm V_MFMA_F32_32X32X4BF16 : MAIInst<"v_mfma_f32_32x32x4bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_32x32x4bf16>;
}
-let SubtargetPredicate = HasGFX950Insts, is_gfx940_xdl = 1 in {
+let SubtargetPredicate = HasGFX950Insts, is_sgemm = 0 in {
defm V_MFMA_F32_16X16X32_F16 : MAIInst<"v_mfma_f32_16x16x32f16", "F32_V8F16_X32", int_amdgcn_mfma_f32_16x16x32_f16>;
defm V_MFMA_F32_32X32X16_F16 : MAIInst<"v_mfma_f32_32x32x16f16", "F32_V8F16_X16", int_amdgcn_mfma_f32_32x32x16_f16>;
defm V_MFMA_F32_16X16X32_BF16 : MAIInst<"v_mfma_f32_16x16x32bf16", "F32_V8BF16_X4", int_amdgcn_mfma_f32_16x16x32_bf16>;
@@ -1054,7 +1054,7 @@ defm V_MFMA_LD_SCALE_B32 : VOP3PInst<"v_mfma_ld_scale_b32", VOP_MFMA_LD_SCALE>;
}
let SubtargetPredicate = isGFX90APlus in {
- let is_gfx940_xdl = 1 in {
+ let is_sgemm = 0 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>;
@@ -1068,17 +1068,17 @@ let SubtargetPredicate = isGFX90APlus in {
}
} // End SubtargetPredicate = isGFX90APlus
-let SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 1 in {
+let SubtargetPredicate = isGFX940Plus, is_sgemm = 0 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>;
-} // End SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 1
+} // End SubtargetPredicate = isGFX940Plus, is_sgemm = 0
-let SubtargetPredicate = HasXF32Insts, is_gfx940_xdl = 1 in {
+let SubtargetPredicate = HasXF32Insts, is_sgemm = 0 in {
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 SubtargetPredicate = HasXF32Insts, is_gfx940_xdl = 1
+} // End SubtargetPredicate = HasXF32Insts, is_sgemm = 0
-let SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 1 in {
+let SubtargetPredicate = HasFP8Insts, is_sgemm = 0 in {
defm V_MFMA_F32_16X16X32_BF8_BF8 : MAIInst<"v_mfma_f32_16x16x32_bf8_bf8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_bf8_bf8>;
defm V_MFMA_F32_16X16X32_BF8_FP8 : MAIInst<"v_mfma_f32_16x16x32_bf8_fp8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_bf8_fp8>;
defm V_MFMA_F32_16X16X32_FP8_BF8 : MAIInst<"v_mfma_f32_16x16x32_fp8_bf8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_fp8_bf8>;
@@ -1087,11 +1087,11 @@ let SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 1 in {
defm V_MFMA_F32_32X32X16_BF8_FP8 : MAIInst<"v_mfma_f32_32x32x16_bf8_fp8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_bf8_fp8>;
defm V_MFMA_F32_32X32X16_FP8_BF8 : MAIInst<"v_mfma_f32_32x32x16_fp8_bf8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_fp8_bf8>;
defm V_MFMA_F32_32X32X16_FP8_FP8 : MAIInst<"v_mfma_f32_32x32x16_fp8_fp8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_fp8_fp8>;
-} // End SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 1
+} // End SubtargetPredicate = HasFP8Insts, is_sgemm = 0
multiclass SMFMACInst<string OpName, string P, SDPatternOperator node> {
let Constraints = "$vdst = $src2", DisableEncoding = "$src2",
- isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1, is_gfx940_xdl = 1 in {
+ isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1, is_sgemm = 0 in {
def _e64 : MAIInst<OpName, !cast<VOPProfileSMFMAC>("VOPProfileSMFMAC_" # P), node>;
}
}
@@ -1105,7 +1105,7 @@ defm V_SMFMAC_I32_16X16X64_I8 : SMFMACInst<"v_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>;
}
-let SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 1 in {
+let SubtargetPredicate = HasFP8Insts, is_sgemm = 0 in {
defm V_SMFMAC_F32_16X16X64_BF8_BF8 : SMFMACInst<"v_smfmac_f32_16x16x64_bf8_bf8", "F32_16X16X64_F8", int_amdgcn_smfmac_f32_16x16x64_bf8_bf8>;
defm V_SMFMAC_F32_16X16X64_BF8_FP8 : SMFMACInst<"v_smfmac_f32_16x16x64_bf8_fp8", "F32_16X16X64_F8", int_amdgcn_smfmac_f32_16x16x64_bf8_fp8>;
defm V_SMFMAC_F32_16X16X64_FP8_BF8 : SMFMACInst<"v_smfmac_f32_16x16x64_fp8_bf8", "F32_16X16X64_F8", int_amdgcn_smfmac_f32_16x16x64_fp8_bf8>;
@@ -1114,7 +1114,7 @@ defm V_SMFMAC_F32_32X32X32_BF8_BF8 : SMFMACInst<"v_smfmac_f32_32x32x32_bf8_bf8",
defm V_SMFMAC_F32_32X32X32_BF8_FP8 : SMFMACInst<"v_smfmac_f32_32x32x32_bf8_fp8", "F32_32X32X32_F8", int_amdgcn_smfmac_f32_32x32x32_bf8_fp8>;
defm V_SMFMAC_F32_32X32X32_FP8_BF8 : SMFMACInst<"v_smfmac_f32_32x32x32_fp8_bf8", "F32_32X32X32_F8", int_amdgcn_smfmac_f32_32x32x32_fp8_bf8>;
defm V_SMFMAC_F32_32X32X32_FP8_FP8 : SMFMACInst<"v_smfmac_f32_32x32x32_fp8_fp8", "F32_32X32X32_F8", int_amdgcn_smfmac_f32_32x32x32_fp8_fp8>;
-} // End SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 1
+} // End SubtargetPredicate = HasFP8Insts, is_sgemm = 0
let SubtargetPredicate = HasGFX950Insts in {
defm V_SMFMAC_F32_16X16X64_F16 : SMFMACInst<"v_smfmac_f32_16x16x64_f16", "F32_16X16X64_F16", int_amdgcn_smfmac_f32_16x16x64_f16>;
@@ -1137,7 +1137,7 @@ def MAIInstInfoTable : GenericTable {
let FilterClass = "MAIInst";
let CppTypeName = "MAIInstInfo";
let Fields = [
- "Opcode", "is_dgemm", "is_gfx940_xdl"
+ "Opcode", "is_dgemm", "is_sgemm"
];
let PrimaryKey = ["Opcode"];
More information about the llvm-commits
mailing list