[AMDGPU][MC][GFX90A] Correct MIMG dst size validation

[dp]

Summary

Correct validator to enable MIMG dst size checks.

See bug 56712 for more information.

[dp]

Removed duplicates.

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

public:
   bool isGFX11() const {
     return AMDGPU::isGFX11(getSTI());
   }
 
   bool isGFX11Plus() const {
     return AMDGPU::isGFX11Plus(getSTI());
   }
 
+  bool isGFX10_AEncoding() const { return AMDGPU::isGFX10_AEncoding(getSTI()); }
+
   bool isGFX10_BEncoding() const {
     return AMDGPU::isGFX10_BEncoding(getSTI());
   }
 
   bool hasInv2PiInlineImm() const {
     return getFeatureBits()[AMDGPU::FeatureInv2PiInlineImm];
   }
 
private:
   bool validateSMEMOffset(const MCInst &Inst, const OperandVector &Operands);
   bool validateSOPLiteral(const MCInst &Inst) const;
   bool validateConstantBusLimitations(const MCInst &Inst, const OperandVector &Operands);
   bool validateEarlyClobberLimitations(const MCInst &Inst, const OperandVector &Operands);
   bool validateIntClampSupported(const MCInst &Inst);
   bool validateMIMGAtomicDMask(const MCInst &Inst);
   bool validateMIMGGatherDMask(const MCInst &Inst);
   bool validateMovrels(const MCInst &Inst, const OperandVector &Operands);
-  Optional<StringRef> validateMIMGDataSize(const MCInst &Inst);
+  bool validateMIMGDataSize(const MCInst &Inst, const SMLoc &IDLoc);
   bool validateMIMGAddrSize(const MCInst &Inst);
   bool validateMIMGD16(const MCInst &Inst);
   bool validateMIMGDim(const MCInst &Inst);
   bool validateMIMGMSAA(const MCInst &Inst);
   bool validateOpSel(const MCInst &Inst);
   bool validateDPP(const MCInst &Inst, const OperandVector &Operands);
   bool validateVccOperand(unsigned Reg) const;
   bool validateVOPLiteral(const MCInst &Inst, const OperandVector &Operands);
if ((Desc.TSFlags & SIInstrFlags::IntClamp) != 0 && !hasIntClamp()) {
     int ClampIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::clamp);
     assert(ClampIdx != -1);
     return Inst.getOperand(ClampIdx).getImm() == 0;
   }
 
   return true;
 }
 
-Optional<StringRef> AMDGPUAsmParser::validateMIMGDataSize(const MCInst &Inst) {
+bool AMDGPUAsmParser::validateMIMGDataSize(const MCInst &Inst,
+                                           const SMLoc &IDLoc) {
 
   const unsigned Opc = Inst.getOpcode();
   const MCInstrDesc &Desc = MII.get(Opc);
 
   if ((Desc.TSFlags & SIInstrFlags::MIMG) == 0)
-    return None;
+    return true;
 
   int VDataIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::vdata);
   int DMaskIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::dmask);
   int TFEIdx   = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::tfe);
 
   assert(VDataIdx != -1);
 
-  if (DMaskIdx == -1 || TFEIdx == -1) // intersect_ray
-    return None;
+  if ((DMaskIdx == -1 || TFEIdx == -1) && isGFX10_AEncoding()) // intersect_ray
+    return true;
 
   unsigned VDataSize = AMDGPU::getRegOperandSize(getMRI(), Desc, VDataIdx);
   unsigned TFESize = (TFEIdx != -1 && Inst.getOperand(TFEIdx).getImm()) ? 1 : 0;
   unsigned DMask = Inst.getOperand(DMaskIdx).getImm() & 0xf;
   if (DMask == 0)
     DMask = 1;
 
-  bool isPackedD16 = false;
+  bool IsPackedD16 = false;
   unsigned DataSize =
     (Desc.TSFlags & SIInstrFlags::Gather4) ? 4 : countPopulation(DMask);
   if (hasPackedD16()) {
     int D16Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::d16);
-    isPackedD16 = D16Idx >= 0;
-    if (isPackedD16 && Inst.getOperand(D16Idx).getImm())
+    IsPackedD16 = D16Idx >= 0;
+    if (IsPackedD16 && Inst.getOperand(D16Idx).getImm())
       DataSize = (DataSize + 1) / 2;
   }
 
   if ((VDataSize / 4) == DataSize + TFESize)
-    return None;
+    return true;
 
-  return StringRef(isPackedD16
-                       ? "image data size does not match dmask, d16 and tfe"
-                       : "image data size does not match dmask and tfe");
+  StringRef Modifiers;
+  if (isGFX90A())
+    Modifiers = IsPackedD16 ? "dmask and d16" : "dmask";
+  else
+    Modifiers = IsPackedD16 ? "dmask, d16 and tfe" : "dmask and tfe";
+
+  Error(IDLoc, Twine("image data size does not match ") + Modifiers);
+  return false;
 }
 
 bool AMDGPUAsmParser::validateMIMGAddrSize(const MCInst &Inst) {
   const unsigned Opc = Inst.getOpcode();
   const MCInstrDesc &Desc = MII.get(Opc);
 
   if ((Desc.TSFlags & SIInstrFlags::MIMG) == 0 || !isGFX10Plus())
     return true;
if (!validateMIMGDim(Inst)) {
     Error(IDLoc, "dim modifier is required on this GPU");
     return false;
   }
   if (!validateMIMGMSAA(Inst)) {
     Error(getImmLoc(AMDGPUOperand::ImmTyDim, Operands),
           "invalid dim; must be MSAA type");
     return false;
   }
-  if (auto ErrMsg = validateMIMGDataSize(Inst)) {
-    Error(IDLoc, *ErrMsg);
+  if (!validateMIMGDataSize(Inst, IDLoc)) {
     return false;
   }
   if (!validateMIMGAddrSize(Inst)) {
     Error(IDLoc,
       "image address size does not match dim and a16");
     return false;
   }
   if (!validateMIMGAtomicDMask(Inst)) {

llvm/test/MC/AMDGPU/mimg-err.s

 // RUN: not llvm-mc -arch=amdgcn %s 2>&1 | FileCheck %s --check-prefix=NOGCN --implicit-check-not=error:
 // RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck %s --check-prefix=NOGCN --implicit-check-not=error:
 // RUN: not llvm-mc -arch=amdgcn -mcpu=fiji %s 2>&1 | FileCheck %s --check-prefix=NOGCN --implicit-check-not=error:
+// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck %s --check-prefix=NOGFX9 --implicit-check-not=error:
+// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck %s --check-prefix=NOGFX90A --implicit-check-not=error:
 
 //===----------------------------------------------------------------------===//
 // Image Load/Store
 //===----------------------------------------------------------------------===//
 
 image_load    v[4:6], v[237:240], s[28:35] dmask:0x7 tfe
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: operands are not valid for this GPU or mode
 
 image_load    v[4:5], v[237:240], s[28:35] dmask:0x7
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: image data size does not match dmask and d16
 
 image_store   v[4:7], v[237:240], s[28:35] dmask:0x7
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: image data size does not match dmask and d16
 
 image_store   v[4:7], v[237:240], s[28:35] dmask:0xe
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: image data size does not match dmask and d16
 
 image_load    v4, v[237:240], s[28:35] tfe
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: operands are not valid for this GPU or mode
 
 //===----------------------------------------------------------------------===//
 // Image Sample
 //===----------------------------------------------------------------------===//
 
 image_sample  v[193:195], v[237:240], s[28:35], s[4:7] dmask:0x7 tfe
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: operands are not valid for this GPU or mode
 
 image_sample  v[193:195], v[237:240], s[28:35], s[4:7] dmask:0x3
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: image data size does not match dmask and d16
 
 image_sample  v[193:195], v[237:240], s[28:35], s[4:7] dmask:0xf
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask, d16 and tfe
+// NOGFX90A: error: image data size does not match dmask and d16
 
 //===----------------------------------------------------------------------===//
 // Image Atomics
 //===----------------------------------------------------------------------===//
 
 image_atomic_add v252, v2, s[8:15] dmask:0x1 tfe
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask and tfe
+// NOGFX90A: error: operands are not valid for this GPU or mode
 
 image_atomic_add v[6:7], v255, s[8:15] dmask:0x2
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask and tfe
+// NOGFX90A: error: image data size does not match dmask
 
 image_atomic_add v[6:7], v255, s[8:15] dmask:0xf
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask and tfe
+// NOGFX90A: error: image data size does not match dmask
 
 image_atomic_cmpswap v[4:7], v[192:195], s[28:35] dmask:0xf tfe
-// NOGCN: error: image data size does not match dmask and tfe
+// NOGCN:    error: image data size does not match dmask and tfe
+// NOGFX9:   error: image data size does not match dmask and tfe
+// NOGFX90A: error: operands are not valid for this GPU or mode
 
 image_atomic_add v252, v2, s[8:15]
-// NOGCN: error: invalid atomic image dmask
+// NOGCN:    error: invalid atomic image dmask
+// NOGFX9:   error: invalid atomic image dmask
+// NOGFX90A: error: invalid atomic image dmask
 
 image_atomic_add v[6:7], v255, s[8:15] dmask:0x2 tfe
-// NOGCN: error: invalid atomic image dmask
+// NOGCN:    error: invalid atomic image dmask
+// NOGFX9:   error: invalid atomic image dmask
+// NOGFX90A: error: operands are not valid for this GPU or mode
 
 image_atomic_cmpswap v[4:7], v[192:195], s[28:35] dmask:0xe tfe
-// NOGCN: error: invalid atomic image dmask
+// NOGCN:    error: invalid atomic image dmask
+// NOGFX9:   error: invalid atomic image dmask
+// NOGFX90A: error: operands are not valid for this GPU or mode
 
 //===----------------------------------------------------------------------===//
 // Image Gather
 //===----------------------------------------------------------------------===//
 
 image_gather4_cl v[5:8], v[1:4], s[8:15], s[12:15] dmask:0x3
-// NOGCN: error: invalid image_gather dmask: only one bit must be set
+// NOGCN:    error: invalid image_gather dmask: only one bit must be set
+// NOGFX9:   error: invalid image_gather dmask: only one bit must be set
+// NOGFX90A: error: instruction not supported on this GPU

llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s

 // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck --check-prefixes=GFX90A --implicit-check-not=error: %s
 
 v_add_f64 v[1:2], v[1:2], v[1:2]
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
 
 global_load_dwordx3 v[1:3], v[0:1], off
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
 
 global_load_dwordx4 v[1:4], v[0:1], off
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
 
-image_load v[1:5], v2, s[0:7] dmask:0xf unorm
+image_load v[1:4], v2, s[0:7] dmask:0xf unorm
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
 
 v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15]
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
 
 v_mfma_i32_4x4x4i8 a[1:4], a0, v1, 2
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
 
 v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[17:32]
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
 
 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[33:64]
 // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned