Index: llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -1448,6 +1448,8 @@
     return AMDGPU::isGFX11Plus(getSTI());
   }
 
+  bool isGFX10_AEncoding() const { return AMDGPU::isGFX10_AEncoding(getSTI()); }
+
   bool isGFX10_BEncoding() const {
     return AMDGPU::isGFX10_BEncoding(getSTI());
   }
@@ -3569,7 +3571,7 @@
 
   assert(VDataIdx != -1);
 
-  if (DMaskIdx == -1 || TFEIdx == -1) // intersect_ray
+  if ((DMaskIdx == -1 || TFEIdx == -1) && isGFX10_AEncoding()) // intersect_ray
     return None;
 
   unsigned VDataSize = AMDGPU::getRegOperandSize(getMRI(), Desc, VDataIdx);
@@ -3591,6 +3593,11 @@
   if ((VDataSize / 4) == DataSize + TFESize)
     return None;
 
+  if (isGFX90A())
+    return StringRef(isPackedD16
+                         ? "image data size does not match dmask and d16"
+                         : "image data size does not match dmask");
+
   return StringRef(isPackedD16
                        ? "image data size does not match dmask, d16 and tfe"
                        : "image data size does not match dmask and tfe");
Index: llvm/test/MC/AMDGPU/mimg-err.s
===================================================================
--- llvm/test/MC/AMDGPU/mimg-err.s
+++ 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
Index: llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
===================================================================
--- llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -9,7 +9,7 @@
 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]