From 071403e67e4971c020ba2d82a901408f5ebe2744 Mon Sep 17 00:00:00 2001 From: Shilei Tian Date: Sun, 9 Aug 2026 15:52:45 -0400 Subject: [PATCH] [AMDGPU][MC] Return fail when it is not a VGPR in the decoder Fixes #215003. --- .../Disassembler/AMDGPUDisassembler.cpp | 8 ++++++ .../Disassembler/AMDGPU/gfx9_dasm_mai_err.txt | 26 +++++++++++++++++++ 2 files changed, 34 insertions(+) create mode 100644 llvm/test/MC/Disassembler/AMDGPU/gfx9_dasm_mai_err.txt diff --git a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp index 90bc97a32a7f2..1e028891e5ad0 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 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 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..1e8c504025138 --- /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=amdgpu9.08-amd-amdhsa -disassemble -show-encoding < %s 2>&1 | FileCheck --implicit-check-not=warning: %s +# RUN: llvm-mc -triple=amdgpu9.0a-amd-amdhsa -disassemble -show-encoding < %s 2>&1 | FileCheck --implicit-check-not=warning: %s +# RUN: llvm-mc -triple=amdgpu9.42-amd-amdhsa -disassemble -show-encoding < %s 2>&1 | FileCheck --implicit-check-not=warning: %s +# RUN: llvm-mc -triple=amdgpu9.50-amd-amdhsa -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