[AMDGPU] Check for immediate SrcC in mfma in AsmParser

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

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@369819 91177308-0d34-0410-b5e6-96231b3b80d8

lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

@@ -3283,6 +3283,10 @@ bool AMDGPUAsmParser::validateVOP3Literal(const MCInst &Inst) const {
     if (!MO.isImm() || !AMDGPU::isSISrcOperand(Desc, OpIdx))
       continue;
 
+    if (OpIdx == Src2Idx && (Desc.TSFlags & SIInstrFlags::IsMAI) &&
+        getFeatureBits()[AMDGPU::FeatureMFMAInlineLiteralBug])
+      return false;
+
     if (!isInlineConstant(Inst, OpIdx)) {
       uint32_t Value = static_cast<uint32_t>(MO.getImm());
       if (NumLiterals == 0 || LiteralValue != Value) {

test/MC/AMDGPU/mai-err.s

@@ -48,3 +48,480 @@ v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 65
 
 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 0
 // GFX900: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0
+// GFX908: error: invalid literal operand
+
+v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
+// GFX908: error: invalid literal operand

test/MC/AMDGPU/mai.s

@@ -48,30 +48,6 @@ v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32]
 v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[1:16]
 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x06,0x04]
 
@@ -96,30 +72,6 @@ v_mfma_f32_16x16x1f32 a[0:15], a0, a1, a[1:16]
 v_mfma_f32_16x16x1f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0 ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0 ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0 ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0 ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[1:4]
 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x06,0x04]
 
@@ -144,30 +96,6 @@ v_mfma_f32_4x4x1f32 a[0:3], a0, a1, a[1:4]
 v_mfma_f32_4x4x1f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0 ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0 ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0 ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0 ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[1:16]
 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x06,0x04]
 
@@ -192,30 +120,6 @@ v_mfma_f32_32x32x2f32 a[0:15], a0, a1, a[1:16]
 v_mfma_f32_32x32x2f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0 ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0 ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0 ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0 ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[1:4]
 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x06,0x04]
 
@@ -240,30 +144,6 @@ v_mfma_f32_16x16x4f32 a[0:3], a0, a1, a[1:4]
 v_mfma_f32_16x16x4f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0 ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0 ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0 ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0 ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], a[1:32]
 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], a[1:32] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0x06,0x04]
 
@@ -288,30 +168,6 @@ v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], a[1:32]
 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], a[1:32] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], a[1:16]
 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], a[1:16] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0x06,0x04]
 
@@ -336,30 +192,6 @@ v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], a[1:16]
 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], a[1:4]
 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], a[1:4] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0x06,0x04]
 
@@ -384,30 +216,6 @@ v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], a[1:4]
 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], a[1:16]
 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], a[1:16] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0x06,0x04]
 
@@ -432,30 +240,6 @@ v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], a[1:16]
 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], a[1:4]
 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], a[1:4] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0x06,0x04]
 
@@ -480,30 +264,6 @@ v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], a[1:4]
 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0 ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0 ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[1:32]
 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x06,0x04]
 
@@ -528,30 +288,6 @@ v_mfma_i32_32x32x4i8 a[0:31], a0, a1, a[1:32]
 v_mfma_i32_32x32x4i8 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2 ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x0a,0x02]
-
-v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x0a,0xe2]
-
-v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2 ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x0a,0x12]
-
-v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x0a,0xf2]
-
-v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2 ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x0a,0x0a]
-
-v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x0a,0xea]
-
-v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2 ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x0a,0x1a]
-
-v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x0a,0xfa]
-
 v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[1:16]
 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x06,0x04]
 
@@ -576,30 +312,6 @@ v_mfma_i32_16x16x4i8 a[0:15], a0, a1, a[1:16]
 v_mfma_i32_16x16x4i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2 ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x0a,0x02]
-
-v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x0a,0xe2]
-
-v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2 ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x0a,0x12]
-
-v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x0a,0xf2]
-
-v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2 ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x0a,0x0a]
-
-v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x0a,0xea]
-
-v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2 ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x0a,0x1a]
-
-v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x0a,0xfa]
-
 v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[1:4]
 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x06,0x04]
 
@@ -624,30 +336,6 @@ v_mfma_i32_4x4x4i8 a[0:3], a0, a1, a[1:4]
 v_mfma_i32_4x4x4i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2 ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x02]
-
-v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x0a,0xe2]
-
-v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2 ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x12]
-
-v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x0a,0xf2]
-
-v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2 ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x0a]
-
-v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x0a,0xea]
-
-v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2 ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x1a]
-
-v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x0a,0xfa]
-
 v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[1:16]
 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x06,0x04]
 
@@ -672,30 +360,6 @@ v_mfma_i32_32x32x8i8 a[0:15], a0, a1, a[1:16]
 v_mfma_i32_32x32x8i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x0a,0x02]
-
-v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x0a,0xe2]
-
-v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x0a,0x12]
-
-v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x0a,0xf2]
-
-v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x0a,0x0a]
-
-v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x0a,0xea]
-
-v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x0a,0x1a]
-
-v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x0a,0xfa]
-
 v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[1:4]
 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x06,0x04]
 
@@ -720,30 +384,6 @@ v_mfma_i32_16x16x16i8 a[0:3], a0, a1, a[1:4]
 v_mfma_i32_16x16x16i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x0a,0x02]
-
-v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x0a,0xe2]
-
-v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x0a,0x12]
-
-v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x0a,0xf2]
-
-v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x0a,0x0a]
-
-v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x0a,0xea]
-
-v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x0a,0x1a]
-
-v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x0a,0xfa]
-
 v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, a[1:32]
 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0x06,0x04]
 
@@ -768,30 +408,6 @@ v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, a[1:32]
 v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0 ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0 ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0 ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0 ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[1:16]
 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0x06,0x04]
 
@@ -816,30 +432,6 @@ v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, a[1:16]
 v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0 ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0 ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0 ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0 ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[1:4]
 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0x06,0x04]
 
@@ -864,30 +456,6 @@ v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, a[1:4]
 v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0 ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0 ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0 ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0 ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[1:16]
 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0x06,0x04]
 
@@ -912,30 +480,6 @@ v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, a[1:16]
 v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0x06,0xfc]
 
-v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0 ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0 ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0 ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0 ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0xd6,0xfb]
-
 v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[1:4]
 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0x06,0x04]
 
@@ -959,27 +503,3 @@ v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, a[1:4]
 
 v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0x06,0xfc]
-
-v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0 ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0xd6,0x03]
-
-v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0xd6,0xe3]
-
-v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0 ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0xd6,0x13]
-
-v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0xd6,0xf3]
-
-v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0 ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0xd6,0x0b]
-
-v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0xd6,0xeb]
-
-v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0 ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0xd6,0x1b]
-
-v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7
-// GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0xd6,0xfb]