[llvm] r369819 - [AMDGPU] Check for immediate SrcC in mfma in AsmParser
Stanislav Mekhanoshin via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 23 15:22:50 PDT 2019
Author: rampitec
Date: Fri Aug 23 15:22:49 2019
New Revision: 369819
URL: http://llvm.org/viewvc/llvm-project?rev=369819&view=rev
Log:
[AMDGPU] Check for immediate SrcC in mfma in AsmParser
Differential Revision: https://reviews.llvm.org/D66674
Modified:
llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
llvm/trunk/test/MC/AMDGPU/mai-err.s
llvm/trunk/test/MC/AMDGPU/mai.s
Modified: llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp?rev=369819&r1=369818&r2=369819&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp Fri Aug 23 15:22:49 2019
@@ -3283,6 +3283,10 @@ bool AMDGPUAsmParser::validateVOP3Litera
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) {
Modified: llvm/trunk/test/MC/AMDGPU/mai-err.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/mai-err.s?rev=369819&r1=369818&r2=369819&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/mai-err.s (original)
+++ llvm/trunk/test/MC/AMDGPU/mai-err.s Fri Aug 23 15:22:49 2019
@@ -48,3 +48,480 @@ v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 6
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
Modified: llvm/trunk/test/MC/AMDGPU/mai.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/mai.s?rev=369819&r1=369818&r2=369819&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/mai.s (original)
+++ llvm/trunk/test/MC/AMDGPU/mai.s Fri Aug 23 15:22:49 2019
@@ -48,30 +48,6 @@ v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a
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
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:
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
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[
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
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
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:
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
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
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[
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[
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[
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[
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,
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,
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
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,
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
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]
More information about the llvm-commits
mailing list