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

Dmitry Preobrazhensky via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Tue May 26 09:11:49 PDT 2020


dp marked 4 inline comments as done.
dp added inline comments.


================
Comment at: llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp:3167-3170
+  if (isSGPR(mc2PseudoReg(Reg), TRI)) {
+    Error(getLoc(), "source operand must be either a VGPR or an inline constant");
+    return false;
+  }
----------------
arsenm wrote:
> This isn't checking the constant values as the error implies
I'll re-title the change and update description


================
Comment at: llvm/test/MC/AMDGPU/mai.s:42
+// 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]
----------------
arsenm wrote:
> Test with inline immediate and non-inline immediate constants?
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.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D80530/new/

https://reviews.llvm.org/D80530





More information about the llvm-commits mailing list