[AMDGPU][MC] Return fail when it is not a VGPR in the decoder

[shiltian]

Fixes #215003.

---

This PR was assisted by AI but I reviewed all the changes.

[shiltian]

• [AMDGPU][MC] Return fail when it is not a VGPR in the decoder #215136 👈
• main

This stack of pull requests is managed by sgh.

[llvmorg-github-actions]

@llvm/pr-subscribers-backend-amdgpu

Author: Shilei Tian (shiltian)

Changes

Fixes #215003.

---

This PR was assisted by AI but I reviewed all the changes.

---

Full diff: https://github.com/llvm/llvm-project/pull/215136.diff

2 Files Affected:

• (modified) llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp (+8)
• (added) llvm/test/MC/Disassembler/AMDGPU/gfx9_dasm_mai_err.txt (+26)

diff --git a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
index d05997cacba0a..bf7502f4aec6d 100644
--- a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
+++ b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
@@ -265,6 +265,10 @@ static DecodeStatus decodeSrcReg9(MCInst &Inst, unsigned Imm,
 template <unsigned OpWidth>
 static DecodeStatus decodeSrcA9(MCInst &Inst, unsigned Imm, uint64_t /* Addr */,
                                 const MCDisassembler *Decoder) {
+  // A clear Imm{8} names an SGPR or an inline constant, which this
+  // register-only operand cannot hold.
+  if (!(Imm & AMDGPU::EncValues::IS_VGPR))
+    return MCDisassembler::Fail;
   return decodeSrcOp(Inst, 9, OpWidth, Imm, Imm | 512, Decoder);
 }
 
@@ -274,6 +278,10 @@ template <unsigned OpWidth>
 static DecodeStatus decodeSrcAV10(MCInst &Inst, unsigned Imm,
                                   uint64_t /* Addr */,
                                   const MCDisassembler *Decoder) {
+  // A clear Imm{8} names an SGPR or an inline constant, which this
+  // register-only operand cannot hold.
+  if (!(Imm & AMDGPU::EncValues::IS_VGPR))
+    return MCDisassembler::Fail;
   return decodeSrcOp(Inst, 10, OpWidth, Imm, Imm, Decoder);
 }
 
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx9_dasm_mai_err.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx9_dasm_mai_err.txt
new file mode 100644
index 0000000000000..34dd986346bef
--- /dev/null
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx9_dasm_mai_err.txt
@@ -0,0 +1,26 @@
+# NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py UTC_ARGS: --version 6
+# RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx908 -disassemble -show-encoding < %s 2>&1 | FileCheck --implicit-check-not=warning: %s
+# RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx90a -disassemble -show-encoding < %s 2>&1 | FileCheck --implicit-check-not=warning: %s
+# RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx942 -disassemble -show-encoding < %s 2>&1 | FileCheck --implicit-check-not=warning: %s
+# RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx950 -disassemble -show-encoding < %s 2>&1 | FileCheck --implicit-check-not=warning: %s
+
+# The AV and AGPR source operands below only accept vector registers, so an
+# encoding naming an SGPR or an inline constant must be rejected. -show-encoding
+# is required: re-encoding the decoded instruction is what used to assert on the
+# non-register operand.
+
+# This is v_mfma_f32_16x16x1_4b_f32 with s0 as src0.
+0x00,0x00,0xc1,0xd3,0x00,0x00,0x02,0x04
+# CHECK: :[[@LINE-1]]:1: warning: invalid instruction encoding
+
+# This is v_mfma_f32_16x16x1_4b_f32 with the inline constant 0 as src1.
+0x00,0x00,0xc1,0xd3,0x00,0x01,0x01,0x04
+# CHECK: :[[@LINE-1]]:1: warning: invalid instruction encoding
+
+# This is v_accvgpr_read_b32 with s0 as src0.
+0x02,0x40,0xd8,0xd3,0x00,0x00,0x00,0x18
+# CHECK: :[[@LINE-1]]:1: warning: invalid instruction encoding
+
+# This is v_accvgpr_read_b32 with the inline constant 0 as src0.
+0x02,0x40,0xd8,0xd3,0x80,0x00,0x00,0x18
+# CHECK: :[[@LINE-1]]:1: warning: invalid instruction encoding