[AMDGPU][MC][GFX908] Corrected src0 of v_accvgpr_write to accept only VGPRs and inline constants.

This change disables use of special SGPR registers like scc, vccz, execz, etc as operands of v_accvgpr_write.

See bug 45414: https://bugs.llvm.org/show_bug.cgi?id=45414

Reviewers: arsenm, rampitec

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

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

@@ -1346,6 +1346,7 @@ private:
   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 @@ bool AMDGPUAsmParser::validateMovrels(const MCInst &Inst) {
   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 @@ bool AMDGPUAsmParser::validateInstruction(const MCInst &Inst,
   if (!validateSMEMOffset(Inst, Operands)) {
     return false;
   }
+  if (!validateMAIAccWrite(Inst)) {
+    return false;
+  }
 
   return true;
 }

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_b32 a2, v1
 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]