diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index effde63adc3d88..299ed1477c5c1c 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 9c8a9c8abf6433..76c9f259309ab1 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 4056a8ac903334..75725c19b35584 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]