[llvm] 2b230d6 - [AMDGPU][MC][GFX90A] Correct MIMG dst size validation
Dmitry Preobrazhensky via llvm-commits
llvm-commits at lists.llvm.org
Thu Jul 28 04:30:28 PDT 2022
Author: Dmitry Preobrazhensky
Date: 2022-07-28T14:30:08+03:00
New Revision: 2b230d69ad44a4ead1071572e78240e9bdc74e0d
URL: https://github.com/llvm/llvm-project/commit/2b230d69ad44a4ead1071572e78240e9bdc74e0d
DIFF: https://github.com/llvm/llvm-project/commit/2b230d69ad44a4ead1071572e78240e9bdc74e0d.diff
LOG: [AMDGPU][MC][GFX90A] Correct MIMG dst size validation
Correct validator to enable MIMG dst size checks.
Differential Revision: https://reviews.llvm.org/D130512
Added:
Modified:
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
llvm/test/MC/AMDGPU/mimg-err.s
llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index effde63adc3d..299ed1477c5c 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -1448,6 +1448,8 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
return AMDGPU::isGFX11Plus(getSTI());
}
+ bool isGFX10_AEncoding() const { return AMDGPU::isGFX10_AEncoding(getSTI()); }
+
bool isGFX10_BEncoding() const {
return AMDGPU::isGFX10_BEncoding(getSTI());
}
@@ -1630,7 +1632,7 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
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);
@@ -3555,13 +3557,14 @@ bool AMDGPUAsmParser::validateIntClampSupported(const MCInst &Inst) {
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);
@@ -3569,8 +3572,8 @@ Optional<StringRef> AMDGPUAsmParser::validateMIMGDataSize(const MCInst &Inst) {
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;
@@ -3578,22 +3581,27 @@ Optional<StringRef> AMDGPUAsmParser::validateMIMGDataSize(const MCInst &Inst) {
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) {
@@ -4577,8 +4585,7 @@ bool AMDGPUAsmParser::validateInstruction(const MCInst &Inst,
"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)) {
diff --git a/llvm/test/MC/AMDGPU/mimg-err.s b/llvm/test/MC/AMDGPU/mimg-err.s
index 9c8a9c8abf64..76c9f259309a 100644
--- a/llvm/test/MC/AMDGPU/mimg-err.s
+++ b/llvm/test/MC/AMDGPU/mimg-err.s
@@ -1,67 +1,101 @@
// 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
diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index 4056a8ac9033..75725c19b355 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -9,7 +9,7 @@ global_load_dwordx3 v[1:3], v[0:1], off
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]
More information about the llvm-commits
mailing list