[PATCH] D121956: [AMDGPU] Disable some MFMA instructions on gfx940

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG4570527e7210: [AMDGPU] Disable some MFMA instructions on gfx940 (authored by rampitec).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D121956/new/

https://reviews.llvm.org/D121956

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


Index: llvm/test/MC/AMDGPU/mai-err-gfx940.s
===================================================================
--- llvm/test/MC/AMDGPU/mai-err-gfx940.s
+++ 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 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
Index: llvm/lib/Target/AMDGPU/VOP3PInstructions.td
===================================================================
--- llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -503,6 +503,8 @@
 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_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_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_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
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D121956.416590.patch
Type: text/x-patch
Size: 3308 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20220318/de56a15b/attachment.bin>