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

[dp]

Summary

v_accvgpr_write may only use VGPRs and inline constants as src0.
This change disables use of special SGPR registers like scc, vccz, execz, etc.

See bug 45414.

[arsenm]

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

3167–3170

This isn't checking the constant values as the error implies

llvm/test/MC/AMDGPU/mai.s

42

Test with inline immediate and non-inline immediate constants?

[dp]

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

3167–3170

I'll re-title the change and update description

llvm/test/MC/AMDGPU/mai.s

42

A test for an inline immediate already exists - see "v_accvgpr_write_b32 a2, -2" above.
I'll add a test for a non-inline immediate.

[rampitec]

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

1349

Can you rename it to validateMAIAccWrite? This is not a load.

[dp]

Renamed validateMAIAccLoad to validateMAIAccWrite.

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

private:
   bool validateMIMGDataSize(const MCInst &Inst);
   bool validateMIMGAddrSize(const MCInst &Inst);
   bool validateMIMGD16(const MCInst &Inst);
   bool validateMIMGDim(const MCInst &Inst);
   bool validateLdsDirect(const MCInst &Inst);
   bool validateOpSel(const MCInst &Inst);
   bool validateVccOperand(unsigned Reg) const;
   bool validateVOP3Literal(const MCInst &Inst) const;
+  bool validateMAIAccLoad(const MCInst &Inst);
   unsigned getConstantBusLimit(unsigned Opcode) const;
   bool usesConstantBus(const MCInst &Inst, unsigned OpIdx);
   bool isInlineConstant(const MCInst &Inst, unsigned OpIdx) const;
   unsigned findImplicitSGPRReadInVOP(const MCInst &Inst) const;
 
   bool isId(const StringRef Id) const;
   bool isId(const AsmToken &Token, const StringRef Id) const;
   bool isToken(const AsmToken::TokenKind Kind) const;
bool AMDGPUAsmParser::validateMovrels(const MCInst &Inst) {
   if (!Src0.isReg())
     return false;
 
   auto Reg = Src0.getReg();
   const MCRegisterInfo *TRI = getContext().getRegisterInfo();
   return !isSGPR(mc2PseudoReg(Reg), TRI);
 }
 
+bool AMDGPUAsmParser::validateMAIAccLoad(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();
   const MCInstrDesc &Desc = MII.get(Opc);
 
   if ((Desc.TSFlags & SIInstrFlags::MIMG) == 0)
     return true;
 
if (!validateMovrels(Inst)) {
     return false;
   }
   if (!validateFlatOffset(Inst, Operands)) {
     return false;
   }
   if (!validateSMEMOffset(Inst, Operands)) {
     return false;
   }
+  if (!validateMAIAccLoad(Inst)) {
+    return false;
+  }
 
   return true;
 }
 
 static std::string AMDGPUMnemonicSpellCheck(StringRef S,
                                             const FeatureBitset &FBS,
                                             unsigned VariantID = 0);
 

llvm/test/MC/AMDGPU/mai.s

-// 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]
 
 v_accvgpr_read_b32 v2, a1
 // GFX908: v_accvgpr_read_b32 v2, a1       ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08]
 
 v_accvgpr_read_b32 v2, a255

 // GFX908: v_accvgpr_write_b32 a2, -2      ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00]
 
 v_accvgpr_write_b32 a2, v1
 // GFX908: v_accvgpr_write_b32 a2, v1      ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00]
 
 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]
 
 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7
 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xe4]
 
 v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32]
 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x14]