diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp --- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp +++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp @@ -1346,6 +1346,7 @@ bool validateOpSel(const MCInst &Inst); bool validateVccOperand(unsigned Reg) const; bool validateVOP3Literal(const MCInst &Inst) const; + bool validateMAIAccWrite(const MCInst &Inst); unsigned getConstantBusLimit(unsigned Opcode) const; bool usesConstantBus(const MCInst &Inst, unsigned OpIdx); bool isInlineConstant(const MCInst &Inst, unsigned OpIdx) const; @@ -3147,6 +3148,30 @@ return !isSGPR(mc2PseudoReg(Reg), TRI); } +bool AMDGPUAsmParser::validateMAIAccWrite(const MCInst &Inst) { + + const unsigned Opc = Inst.getOpcode(); + + if (Opc != AMDGPU::V_ACCVGPR_WRITE_B32_vi) + return true; + + const int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0); + assert(Src0Idx != -1); + + const MCOperand &Src0 = Inst.getOperand(Src0Idx); + if (!Src0.isReg()) + return true; + + auto Reg = Src0.getReg(); + const MCRegisterInfo *TRI = getContext().getRegisterInfo(); + if (isSGPR(mc2PseudoReg(Reg), TRI)) { + Error(getLoc(), "source operand must be either a VGPR or an inline constant"); + return false; + } + + return true; +} + bool AMDGPUAsmParser::validateMIMGD16(const MCInst &Inst) { const unsigned Opc = Inst.getOpcode(); @@ -3617,6 +3642,9 @@ if (!validateSMEMOffset(Inst, Operands)) { return false; } + if (!validateMAIAccWrite(Inst)) { + return false; + } return true; } diff --git a/llvm/test/MC/AMDGPU/mai.s b/llvm/test/MC/AMDGPU/mai.s --- a/llvm/test/MC/AMDGPU/mai.s +++ b/llvm/test/MC/AMDGPU/mai.s @@ -1,4 +1,5 @@ -// RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s +// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s +// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s 2>&1 | FileCheck -check-prefix=NOGFX908 %s v_accvgpr_read_b32 v2, a0 // GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08] @@ -24,6 +25,24 @@ v_accvgpr_write a2, v255 // GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x00,0xd9,0xd3,0xff,0x01,0x00,0x00] +v_accvgpr_write a2, 100 +// NOGFX908: error: invalid operand for instruction + +v_accvgpr_write a2, execz +// NOGFX908: error: source operand must be either a VGPR or an inline constant + +v_accvgpr_write a2, vccz +// NOGFX908: error: source operand must be either a VGPR or an inline constant + +v_accvgpr_write a2, scc +// NOGFX908: error: source operand must be either a VGPR or an inline constant + +v_accvgpr_write a2, shared_base +// NOGFX908: error: source operand must be either a VGPR or an inline constant + +v_accvgpr_write a2, pops_exiting_wave_id +// NOGFX908: error: source operand must be either a VGPR or an inline constant + v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] // GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x04]