[llvm] 4570527 - [AMDGPU] Disable some MFMA instructions on gfx940

[Stanislav Mekhanoshin via llvm-commits]

Author: Stanislav Mekhanoshin
Date: 2022-03-18T13:19:12-07:00
New Revision: 4570527e7210b4f379f20af36ba4026ddafd852f

URL: https://github.com/llvm/llvm-project/commit/4570527e7210b4f379f20af36ba4026ddafd852f
DIFF: https://github.com/llvm/llvm-project/commit/4570527e7210b4f379f20af36ba4026ddafd852f.diff

LOG: [AMDGPU] Disable some MFMA instructions on gfx940

Differential Revision: https://reviews.llvm.org/D121956

Added: 
    

Modified: 
    llvm/lib/Target/AMDGPU/VOP3PInstructions.td
    llvm/test/MC/AMDGPU/mai-err-gfx940.s

Removed: 
    


################################################################################
diff  --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
index 964c020ab8a82..2f87a75c9496b 100644
--- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -503,6 +503,8 @@ defm V_MFMA_F32_32X32X2F32  : MAIInst<"v_mfma_f32_32x32x2f32",  "F32_F32_X16",
 defm V_MFMA_F32_32X32X4F16  : MAIInst<"v_mfma_f32_32x32x4f16",  "F32_V4F16_X32", int_amdgcn_mfma_f32_32x32x4f16>;
 defm V_MFMA_F32_32X32X8F16  : MAIInst<"v_mfma_f32_32x32x8f16",  "F32_V4F16_X16", int_amdgcn_mfma_f32_32x32x8f16>;
 defm V_MFMA_I32_32X32X4I8   : MAIInst<"v_mfma_i32_32x32x4i8",   "I32_I32_X32",   int_amdgcn_mfma_i32_32x32x4i8>;
+
+let Predicates = [isGFX908orGFX90A] in {
 defm V_MFMA_I32_16X16X16I8  : MAIInst<"v_mfma_i32_16x16x16i8",  "I32_I32_X4",    int_amdgcn_mfma_i32_16x16x16i8>;
 defm V_MFMA_I32_32X32X8I8   : MAIInst<"v_mfma_i32_32x32x8i8",   "I32_I32_X16",   int_amdgcn_mfma_i32_32x32x8i8>;
 defm V_MFMA_F32_4X4X2BF16   : MAIInst<"v_mfma_f32_4x4x2bf16",   "F32_V2I16_X4",  int_amdgcn_mfma_f32_4x4x2bf16>;
@@ -510,6 +512,7 @@ defm V_MFMA_F32_16X16X2BF16 : MAIInst<"v_mfma_f32_16x16x2bf16", "F32_V2I16_X16",
 defm V_MFMA_F32_16X16X8BF16 : MAIInst<"v_mfma_f32_16x16x8bf16", "F32_V2I16_X4",  int_amdgcn_mfma_f32_16x16x8bf16>;
 defm V_MFMA_F32_32X32X2BF16 : MAIInst<"v_mfma_f32_32x32x2bf16", "F32_V2I16_X32", int_amdgcn_mfma_f32_32x32x2bf16>;
 defm V_MFMA_F32_32X32X4BF16 : MAIInst<"v_mfma_f32_32x32x4bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_32x32x4bf16>;
+}
 
 } // End SubtargetPredicate = HasMAIInsts
 
@@ -704,6 +707,7 @@ 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">;
 
+let SubtargetPredicate = isGFX908orGFX90A in {
 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>;
@@ -711,6 +715,7 @@ defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MFMA <0x69>;
 defm V_MFMA_F32_4X4X2BF16   : VOP3P_Real_MFMA <0x6b>;
 defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA <0x6c>;
 defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA <0x6d>;
+}
 
 } // End SubtargetPredicate = HasMAIInsts
 

diff  --git a/llvm/test/MC/AMDGPU/mai-err-gfx940.s b/llvm/test/MC/AMDGPU/mai-err-gfx940.s
index a2832eec76956..499bd691cf3e5 100644
--- a/llvm/test/MC/AMDGPU/mai-err-gfx940.s
+++ b/llvm/test/MC/AMDGPU/mai-err-gfx940.s
@@ -1,5 +1,20 @@
 // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx940 %s 2>&1 | FileCheck -check-prefix=GFX940 %s
 
+v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, 0
+// GFX940: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, 0
+// GFX940: error: instruction not supported on this GPU
+
+v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, 0
+// GFX940: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, 0
+// GFX940: error: operands are not valid for this GPU or mode
+
+v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, 0
+// GFX940: error: instruction not supported on this GPU
+
 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31] neg:[1,0,0]
 // GFX940: error: invalid modifier: neg is not supported
 
@@ -20,3 +35,9 @@ v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] blgp:7
 
 v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] blgp:7
 // GFX940: error: invalid modifier: blgp is not supported
+
+v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[0:15]
+// GFX940: error: instruction not supported on this GPU
+
+v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[0:3]
+// GFX940: error: instruction not supported on this GPU