[llvm] 4467327 - [AMDGPU][MC][GFX940] Add SMFMAC aliases
Dmitry Preobrazhensky via llvm-commits
llvm-commits at lists.llvm.org
Thu May 19 03:43:21 PDT 2022
Author: Dmitry Preobrazhensky
Date: 2022-05-19T13:40:48+03:00
New Revision: 44673278e029d7a6f56c8a3177247026b831720f
URL: https://github.com/llvm/llvm-project/commit/44673278e029d7a6f56c8a3177247026b831720f
DIFF: https://github.com/llvm/llvm-project/commit/44673278e029d7a6f56c8a3177247026b831720f.diff
LOG: [AMDGPU][MC][GFX940] Add SMFMAC aliases
Differential Revision: https://reviews.llvm.org/D125888
Added:
Modified:
llvm/lib/Target/AMDGPU/VOP3PInstructions.td
llvm/test/MC/AMDGPU/mai-gfx940.s
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
index 1c670658f06fa..aecd40f411ec1 100644
--- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -700,12 +700,13 @@ multiclass VOP3P_Real_MFMA<bits<7> op, string GFX940Name = !cast<VOP3_Pseudo>(NA
}
}
-multiclass VOP3P_Real_SMFMAC<bits<7> op> {
+multiclass VOP3P_Real_SMFMAC<bits<7> op, string alias> {
def _gfx940 : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>,
VOP3Pe_SMFMAC <op> {
let AssemblerPredicate = isGFX940Plus;
let DecoderNamespace = "GFX8";
}
+ def : MnemonicAlias<alias, !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic>;
}
defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>;
@@ -822,12 +823,12 @@ defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx940 <0x61, "v_mfma_f32_16x
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">;
-defm V_SMFMAC_F32_16X16X32_F16 : VOP3P_Real_SMFMAC <0x62>;
-defm V_SMFMAC_F32_32X32X16_F16 : VOP3P_Real_SMFMAC <0x64>;
-defm V_SMFMAC_F32_16X16X32_BF16 : VOP3P_Real_SMFMAC <0x66>;
-defm V_SMFMAC_F32_32X32X16_BF16 : VOP3P_Real_SMFMAC <0x68>;
-defm V_SMFMAC_I32_16X16X64_I8 : VOP3P_Real_SMFMAC <0x6a>;
-defm V_SMFMAC_I32_32X32X32_I8 : VOP3P_Real_SMFMAC <0x6c>;
+defm V_SMFMAC_F32_16X16X32_F16 : VOP3P_Real_SMFMAC <0x62, "v_smfmac_f32_16x16x32f16">;
+defm V_SMFMAC_F32_32X32X16_F16 : VOP3P_Real_SMFMAC <0x64, "v_smfmac_f32_32x32x16f16">;
+defm V_SMFMAC_F32_16X16X32_BF16 : VOP3P_Real_SMFMAC <0x66, "v_smfmac_f32_16x16x32bf16">;
+defm V_SMFMAC_F32_32X32X16_BF16 : VOP3P_Real_SMFMAC <0x68, "v_smfmac_f32_32x32x16bf16">;
+defm V_SMFMAC_I32_16X16X64_I8 : VOP3P_Real_SMFMAC <0x6a, "v_smfmac_i32_16x16x64i8">;
+defm V_SMFMAC_I32_32X32X32_I8 : VOP3P_Real_SMFMAC <0x6c, "v_smfmac_i32_32x32x32i8">;
let SubtargetPredicate = HasPackedFP32Ops in {
defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>;
diff --git a/llvm/test/MC/AMDGPU/mai-gfx940.s b/llvm/test/MC/AMDGPU/mai-gfx940.s
index f37733c6b901a..52b035fda971b 100644
--- a/llvm/test/MC/AMDGPU/mai-gfx940.s
+++ b/llvm/test/MC/AMDGPU/mai-gfx940.s
@@ -1,9 +1,17 @@
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx940 -show-encoding %s | FileCheck -check-prefix=GFX940 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefix=GFX90A %s
+//===----------------------------------------------------------------------===//
+// Misc opcodes.
+//===----------------------------------------------------------------------===//
+
v_accvgpr_write_b32 a10, s20
// GFX940: v_accvgpr_write_b32 a10, s20 ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18]
+//===----------------------------------------------------------------------===//
+// MFMA opcodes.
+//===----------------------------------------------------------------------===//
+
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
@@ -460,6 +468,10 @@ v_mfma_f32_32x32x4xf32 a[0:15], v[2:3], v[4:5], a[18:33]
// GFX940: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xbf,0xd3,0x02,0x09,0x4a,0x04]
// GFX90A: error: instruction not supported on this GPU
+//===----------------------------------------------------------------------===//
+// SMFMAC opcodes.
+//===----------------------------------------------------------------------===//
+
v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1
// GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe2,0xd3,0x02,0x09,0x02,0x0c]
// GFX90A: error: instruction not supported on this GPU
@@ -507,3 +519,31 @@ v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1
v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11
// GFX940: v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 ; encoding: [0x0a,0x80,0xec,0xd3,0x02,0x09,0x2e,0x14]
// GFX90A: error: instruction not supported on this GPU
+
+//===----------------------------------------------------------------------===//
+// SMFMAC aliases.
+//===----------------------------------------------------------------------===//
+
+v_smfmac_f32_16x16x32f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1
+// GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe2,0xd3,0x02,0x09,0x02,0x0c]
+// GFX90A: error: instruction not supported on this GPU
+
+v_smfmac_f32_32x32x16f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1
+// GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe4,0xd3,0x02,0x09,0x0a,0x0c]
+// GFX90A: error: instruction not supported on this GPU
+
+v_smfmac_f32_16x16x32bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1
+// GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe6,0xd3,0x02,0x09,0x12,0x0c]
+// GFX90A: error: instruction not supported on this GPU
+
+v_smfmac_f32_32x32x16bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1
+// GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe8,0xd3,0x02,0x09,0x1a,0x0c]
+// GFX90A: error: instruction not supported on this GPU
+
+v_smfmac_i32_16x16x64i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1
+// GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xea,0xd3,0x02,0x09,0x22,0x0c]
+// GFX90A: error: instruction not supported on this GPU
+
+v_smfmac_i32_32x32x32i8 a[10:25], v[2:3], a[4:7], v11
+// GFX940: v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 ; encoding: [0x0a,0x80,0xec,0xd3,0x02,0x09,0x2e,0x14]
+// GFX90A: error: instruction not supported on this GPU
More information about the llvm-commits
mailing list