[PATCH] D80530: [AMDGPU][MC][GFX908] Disabled SGPRs as src0 of v_accvgpr_write

[Dmitry Preobrazhensky via Phabricator via llvm-commits]

dp created this revision.
dp added reviewers: rampitec, arsenm.
Herald added subscribers: llvm-commits, kerbowa, hiraditya, t-tye, tpr, dstuttard, yaxunl, nhaehnle, wdng, jvesely, kzhuravl.
Herald added a project: LLVM.

v_accvgpr_write may only use VGPRs and inline constants as src0.
See bug 45414 <https://bugs.llvm.org/show_bug.cgi?id=45414>.


https://reviews.llvm.org/D80530

Files:
  llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
  llvm/test/MC/AMDGPU/mai.s


Index: llvm/test/MC/AMDGPU/mai.s
===================================================================
--- llvm/test/MC/AMDGPU/mai.s
+++ 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,21 @@
 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, 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]
 
Index: llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ 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 validateMAIAccLoad(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::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();
@@ -3617,6 +3642,9 @@
   if (!validateSMEMOffset(Inst, Operands)) {
     return false;
   }
+  if (!validateMAIAccLoad(Inst)) {
+    return false;
+  }
 
   return true;
 }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D80530.266063.patch
Type: text/x-patch
Size: 3114 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20200525/f4e1f34e/attachment-0001.bin>