[llvm] d9ac55f - [AMDGPU] New MFMA names for existing instructions
Stanislav Mekhanoshin via llvm-commits
llvm-commits at lists.llvm.org
Thu Mar 17 13:35:54 PDT 2022
Author: Stanislav Mekhanoshin
Date: 2022-03-17T13:05:36-07:00
New Revision: d9ac55fab2ec49b11c5ab351b3551a7c7bcb89f0
URL: https://github.com/llvm/llvm-project/commit/d9ac55fab2ec49b11c5ab351b3551a7c7bcb89f0
DIFF: https://github.com/llvm/llvm-project/commit/d9ac55fab2ec49b11c5ab351b3551a7c7bcb89f0.diff
LOG: [AMDGPU] New MFMA names for existing instructions
Old names are supported as aliases.
_1k MFMA got new opcodes.
Differential Revision: https://reviews.llvm.org/D121741
Added:
Modified:
llvm/lib/Target/AMDGPU/VOP3PInstructions.td
llvm/test/MC/AMDGPU/mai-gfx940.s
llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
index bcf4ae3ccea6a..964c020ab8a82 100644
--- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -573,16 +573,56 @@ multiclass VOP3P_Real_MFMA_gfx90a<bits<7> op> {
VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64").Pfl, 0>;
} // End AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A"
}
+}
+
+multiclass VOP3P_Real_MFMA_gfx940_aliases<string NameFrom, string NameTo, string Op,
+ VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(Op # "_e64"),
+ VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(Op # "_vgprcd" # "_e64"),
+ VOPProfile Pfl_ACD = PS_ACD.Pfl,
+ VOPProfile Pfl_VCD = PS_VCD.Pfl> {
+ let Predicates = [isGFX940Plus] in {
+ foreach _ = BoolToList<!ne(NameFrom, NameTo)>.ret in {
+ def : InstAlias <NameTo # " " # PS_ACD.AsmOperands,
+ (!cast<VOP3P_Real>(Op # "_gfx940_acd") Pfl_ACD.DstRC:$vdst,
+ Pfl_ACD.Src0RC64:$src0, Pfl_ACD.Src1RC64:$src1, Pfl_ACD.Src2RC64:$src2,
+ cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl;
+ def : InstAlias <NameTo # " " # PS_VCD.AsmOperands,
+ (!cast<VOP3P_Real>(Op # "_gfx940_vcd") Pfl_VCD.DstRC:$vdst,
+ Pfl_VCD.Src0RC64:$src0, Pfl_VCD.Src1RC64:$src1, Pfl_VCD.Src2RC64:$src2,
+ cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl;
+ }
+ } // End Predicates = [isGFX940Plus]
+}
+
+multiclass VOP3P_Real_MFMA_gfx940<bits<7> op, string Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic,
+ VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(NAME # "_e64"),
+ VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64")> {
+ let SubtargetPredicate = isGFX940Plus,
+ AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9",
+ AsmString = Name # PS_ACD.AsmOperands, Constraints = "" in {
+ def _gfx940_acd : VOP3P_Real<PS_ACD, SIEncodingFamily.GFX940>,
+ VOP3Pe_MAI <op, PS_ACD.Pfl, 1>;
+
+ def _gfx940_vcd : VOP3P_Real<PS_VCD, SIEncodingFamily.GFX940>,
+ VOP3Pe_MAI <op, PS_VCD.Pfl, 0>;
+ } // End AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9"
+
+ defm : VOP3P_Real_MFMA_gfx940_aliases<Name, PS_ACD.Mnemonic, NAME>;
-multiclass VOP3P_Real_MFMA<bits<7> op> :
- VOP3P_Real_MFMA_gfx90a <op> {
+ foreach _ = BoolToList<!ne(!subst("_1k", "", PS_ACD.Mnemonic), PS_ACD.Mnemonic)>.ret in
+ defm : VOP3P_Real_MFMA_gfx940_aliases<Name, !subst("_1k", "", PS_ACD.Mnemonic), NAME>;
+}
+
+multiclass VOP3P_Real_MFMA<bits<7> op, string GFX940Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic> :
+ VOP3P_Real_MFMA_gfx90a <op>,
+ VOP3P_Real_MFMA_gfx940 <op, GFX940Name> {
def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>,
VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, ?> {
let AssemblerPredicate = HasMAIInsts;
let DecoderNamespace = "GFX8";
+ let Constraints = "";
}
}
-}
defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>;
defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>;
@@ -650,19 +690,20 @@ let SubtargetPredicate = HasMAIInsts in {
defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>;
defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>;
-defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40>;
-defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41>;
-defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42>;
-defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44>;
-defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45>;
-defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48>;
-defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49>;
-defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a>;
-defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c>;
-defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d>;
-defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50>;
-defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51>;
-defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52>;
+defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40, "v_mfma_f32_32x32x1_2b_f32">;
+defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41, "v_mfma_f32_16x16x1_4b_f32">;
+defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42, "v_mfma_f32_4x4x1_16b_f32">;
+defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44, "v_mfma_f32_32x32x2_f32">;
+defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45, "v_mfma_f32_16x16x4_f32">;
+defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48, "v_mfma_f32_32x32x4_2b_f16">;
+defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49, "v_mfma_f32_16x16x4_4b_f16">;
+defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a, "v_mfma_f32_4x4x4_16b_f16">;
+defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c, "v_mfma_f32_32x32x8_f16">;
+defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d, "v_mfma_f32_16x16x16_f16">;
+defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50, "v_mfma_i32_32x32x4_2b_i8">;
+defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51, "v_mfma_i32_16x16x4_4b_i8">;
+defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52, "v_mfma_i32_4x4x4_16b_i8">;
+
defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MFMA <0x55>;
defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MFMA <0x54>;
defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MFMA <0x68>;
@@ -681,6 +722,15 @@ defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx90a <0x67>;
defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx90a <0x6e>;
defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx90a <0x6f>;
+defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5d, "v_mfma_f32_32x32x4_2b_bf16">;
+defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5e, "v_mfma_f32_16x16x4_4b_bf16">;
+defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5f, "v_mfma_f32_4x4x4_16b_bf16">;
+defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx940 <0x60, "v_mfma_f32_32x32x8_bf16">;
+defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx940 <0x61, "v_mfma_f32_16x16x16_bf16">;
+
+defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx940 <0x6e, "v_mfma_f64_16x16x4_f64">;
+defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx940 <0x6f, "v_mfma_f64_4x4x4_4b_f64">;
+
let SubtargetPredicate = HasPackedFP32Ops in {
defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>;
defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>;
diff --git a/llvm/test/MC/AMDGPU/mai-gfx940.s b/llvm/test/MC/AMDGPU/mai-gfx940.s
index fce8c6a5a65c5..1c58a7a304db2 100644
--- a/llvm/test/MC/AMDGPU/mai-gfx940.s
+++ b/llvm/test/MC/AMDGPU/mai-gfx940.s
@@ -4,3 +4,339 @@
v_accvgpr_write_b32 a10, s20
// GFX940: v_accvgpr_write_b32 a10, s20 ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18]
// GFX90A: error: source operand must be either a VGPR or an inline constant
+
+v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3]
+// GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3]
+// GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f64_4x4x4f64 a[0:1], v[0:1], a[2:3], a[2:3]
+// GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14]
+
+v_mfma_f64_4x4x4f64 v[0:1], v[0:1], a[2:3], v[2:3]
+// GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14]
+
+v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]
+// GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7]
+// GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]
+// GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04]
+
+v_mfma_f64_16x16x4f64 v[0:7], v[0:1], v[2:3], v[0:7]
+// GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04]
+
+v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33]
+// GFX940: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33]
+// GFX940: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[18:33]
+// GFX940: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04]
+
+v_mfma_f32_16x16x1f32 v[0:15], v0, v1, v[18:33]
+// GFX940: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04]
+
+v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5]
+// GFX940: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]
+// GFX940: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[2:5]
+// GFX940: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04]
+
+v_mfma_f32_4x4x1f32 v[0:3], v0, v1, v[2:5]
+// GFX940: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04]
+
+v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33]
+// GFX940: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33]
+// GFX940: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[18:33]
+// GFX940: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04]
+
+v_mfma_f32_32x32x2f32 v[0:15], v0, v1, v[18:33]
+// GFX940: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04]
+
+v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5]
+// GFX940: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5]
+// GFX940: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[2:5]
+// GFX940: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04]
+
+v_mfma_f32_16x16x4f32 v[0:3], v0, v1, v[2:5]
+// GFX940: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04]
+
+v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[2:3], a[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04]
+
+v_mfma_f32_32x32x4f16 v[0:31], v[0:1], v[2:3], v[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04]
+
+v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33]
+// GFX940: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33]
+// GFX940: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[2:3], a[18:33]
+// GFX940: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04]
+
+v_mfma_f32_16x16x4f16 v[0:15], v[0:1], v[2:3], v[18:33]
+// GFX940: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04]
+
+v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[2:3], a[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04]
+
+v_mfma_f32_4x4x4f16 v[0:3], v[0:1], v[2:3], v[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04]
+
+v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33]
+// GFX940: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33]
+// GFX940: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[2:3], a[18:33]
+// GFX940: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04]
+
+v_mfma_f32_32x32x8f16 v[0:15], v[0:1], v[2:3], v[18:33]
+// GFX940: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04]
+
+v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5]
+// GFX940: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5]
+// GFX940: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[2:3], a[2:5]
+// GFX940: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04]
+
+v_mfma_f32_16x16x16f16 v[0:3], v[0:1], v[2:3], v[2:5]
+// GFX940: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04]
+
+v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65]
+// GFX940: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65]
+// GFX940: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[34:65]
+// GFX940: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04]
+
+v_mfma_i32_32x32x4i8 v[0:31], v0, a1, v[34:65]
+// GFX940: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14]
+
+v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33]
+// GFX940: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33]
+// GFX940: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[18:33]
+// GFX940: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04]
+
+v_mfma_i32_16x16x4i8 v[0:15], v0, v1, v[18:33]
+// GFX940: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04]
+
+v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5]
+// GFX940: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5]
+// GFX940: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[2:5]
+// GFX940: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04]
+
+v_mfma_i32_4x4x4i8 v[0:3], v0, v1, v[2:5]
+// GFX940: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04]
+
+v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7
+// GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7
+// GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7
+// GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7
+// GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[34:65] blgp:7
+// GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
+
+v_mfma_f32_32x32x1f32 v[0:31], v0, v1, v[34:65] blgp:7
+// GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
+
+v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65]
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
+// GFX90A: error: operands are not valid for this GPU or mode
+
+v_mfma_f32_32x32x4bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
+// GFX90A: error: operands are not valid for this GPU or mode
+
+v_mfma_f32_32x32x4bf16_1k v[0:31], v[2:3], v[4:5], v[34:65] blgp:5
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
+
+v_mfma_f32_32x32x4bf16_1k a[0:31], v[2:3], v[4:5], a[34:65] blgp:5
+// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
+
+v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33]
+// GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33]
+// GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5
+// GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5
+// GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x4bf16_1k v[0:15], v[2:3], v[4:5], v[18:33] blgp:5
+// GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4]
+
+v_mfma_f32_16x16x4bf16_1k a[0:15], v[2:3], v[4:5], a[18:33] blgp:5
+// GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4]
+
+v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x4bf16 v[0:3], v[2:3], v[4:5], v[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x4bf16 a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x4bf16_1k v[0:3], v[2:3], v[4:5], v[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+
+v_mfma_f32_4x4x4bf16_1k a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+
+v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33]
+// GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33]
+// GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x8bf16 v[0:15], v[2:3], v[4:5], v[18:33]
+// GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x8bf16 a[0:15], v[2:3], v[4:5], a[18:33]
+// GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x8bf16_1k v[0:15], v[2:3], v[4:5], v[18:33]
+// GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]
+
+v_mfma_f32_32x32x8bf16_1k a[0:15], v[2:3], v[4:5], a[18:33]
+// GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]
+
+v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5]
+// GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x16bf16 v[0:3], v[2:3], v[4:5], v[2:5]
+// GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x16bf16 a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]
+// GFX90A: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x16bf16_1k v[0:3], v[2:3], v[4:5], v[2:5]
+// GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]
+
+v_mfma_f32_16x16x16bf16_1k a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt b/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt
index 6dd6ee936d3ba..9c0ff56743f87 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt
@@ -2,3 +2,33 @@
# GFX940: v_accvgpr_write_b32 a10, s20 ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18]
0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18
+
+# GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[2:33] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x00,0xdd,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[2:33] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x80,0xdd,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[2:17] ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x00,0xde,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[2:17] ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x80,0xde,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[2:17] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x00,0xe0,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[2:17] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x80,0xe0,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04
+
+# GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]
+0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04
More information about the llvm-commits
mailing list