diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index c2aa181fb28fb..2139884d000ba 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -267,6 +267,15 @@ void Sema::Initialize() {
   if (getLangOpts().SYCLIsDevice) {
     addImplicitTypedef("__ocl_event_t", Context.OCLEventTy);
     addImplicitTypedef("__ocl_sampler_t", Context.OCLSamplerTy);
+#ifdef SEMA_STRINGIZE
+#error "Undefine SEMA_STRINGIZE macro."
+#endif
+#define SEMA_STRINGIZE(s) #s
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  addImplicitTypedef(SEMA_STRINGIZE(__ocl_##ImgType##_##Suffix##_t),           \
+                     Context.SingletonId);
+#include "clang/Basic/OpenCLImageTypes.def"
+#undef SEMA_STRINGIZE
   }
 
   // Initialize predefined OpenCL types and supported extensions and (optional)
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 552cedcafd279..c0d39f945671b 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -135,11 +135,86 @@ class accessor {
   template <typename... T>
   void use(T... args) const {}
   _ImplT<dimensions> impl;
+
 private:
   void __init(__global dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
 };
 
+template <int dimensions, access::mode accessmode, access::target accesstarget>
+struct opencl_image_type;
+
+#define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \
+  template <>                                                       \
+  struct opencl_image_type<dim, access::mode::accessmode,           \
+                           access::target::Target> {                \
+    using type = __ocl_image##dim##d_##ifarray_##amsuffix##_t;      \
+  };
+
+#define IMAGETY_READ_3_DIM_IMAGE       \
+  IMAGETY_DEFINE(1, read, ro, image, ) \
+  IMAGETY_DEFINE(2, read, ro, image, ) \
+  IMAGETY_DEFINE(3, read, ro, image, )
+
+#define IMAGETY_WRITE_3_DIM_IMAGE       \
+  IMAGETY_DEFINE(1, write, wo, image, ) \
+  IMAGETY_DEFINE(2, write, wo, image, ) \
+  IMAGETY_DEFINE(3, write, wo, image, )
+
+#define IMAGETY_READ_2_DIM_IARRAY                  \
+  IMAGETY_DEFINE(1, read, ro, image_array, array_) \
+  IMAGETY_DEFINE(2, read, ro, image_array, array_)
+
+#define IMAGETY_WRITE_2_DIM_IARRAY                  \
+  IMAGETY_DEFINE(1, write, wo, image_array, array_) \
+  IMAGETY_DEFINE(2, write, wo, image_array, array_)
+
+IMAGETY_READ_3_DIM_IMAGE
+IMAGETY_WRITE_3_DIM_IMAGE
+
+IMAGETY_READ_2_DIM_IARRAY
+IMAGETY_WRITE_2_DIM_IARRAY
+
+template <int dim, access::mode accessmode, access::target accesstarget>
+struct _ImageImplT {
+#ifdef __SYCL_DEVICE_ONLY__
+  typename opencl_image_type<dim, accessmode, accesstarget>::type MImageObj;
+#else
+  range<dim> AccessRange;
+  range<dim> MemRange;
+  id<dim> Offset;
+#endif
+};
+
+template <typename dataT, int dimensions, access::mode accessmode>
+class accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
+public:
+  void use(void) const {}
+  template <typename... T>
+  void use(T... args) {}
+  template <typename... T>
+  void use(T... args) const {}
+  _ImageImplT<dimensions, accessmode, access::target::image> impl;
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(typename opencl_image_type<dimensions, accessmode, access::target::image>::type ImageObj) { impl.MImageObj = ImageObj; }
+#endif
+};
+
+template <typename dataT, int dimensions, access::mode accessmode>
+class accessor<dataT, dimensions, accessmode, access::target::host_image, access::placeholder::false_t> {
+public:
+  void use(void) const {}
+  template <typename... T>
+  void use(T... args) {}
+  template <typename... T>
+  void use(T... args) const {}
+  _ImageImplT<dimensions, accessmode, access::target::host_image> impl;
+};
+
+// TODO: Add support for image_array accessor.
+// template <typename dataT, int dimensions, access::mode accessmode>
+//class accessor<dataT, dimensions, accessmode, access::target::image_array, access::placeholder::false_t>
+
 class kernel {};
 class context {};
 class device {};
@@ -241,13 +316,81 @@ class buffer {
   accessor<T, dimensions, mode, access::target::host_buffer,
            access::placeholder::false_t>
   get_access() {
-    accessor<T, dimensions, mode, access::target::host_buffer,
-             access::placeholder::false_t>{};
+    return accessor<T, dimensions, mode, access::target::host_buffer,
+                    access::placeholder::false_t>{};
   }
 
   template <typename Destination>
   void set_final_data(Destination finalData = nullptr) {}
 };
 
+enum class image_channel_order : unsigned int {
+  a,
+  r,
+  rx,
+  rg,
+  rgx,
+  ra,
+  rgb,
+  rgbx,
+  rgba,
+  argb,
+  bgra,
+  intensity,
+  luminance,
+  abgr
+};
+
+enum class image_channel_type : unsigned int {
+  snorm_int8,
+  snorm_int16,
+  unorm_int8,
+  unorm_int16,
+  unorm_short_565,
+  unorm_short_555,
+  unorm_int_101010,
+  signed_int8,
+  signed_int16,
+  signed_int32,
+  unsigned_int8,
+  unsigned_int16,
+  unsigned_int32,
+  fp16,
+  fp32
+};
+
+template <int dimensions = 1, typename AllocatorT = int>
+class image {
+public:
+  image(image_channel_order Order, image_channel_type Type,
+        const range<dimensions> &Range, const property_list &PropList = {}) {}
+
+  /* -- common interface members -- */
+
+  image(const image &rhs) = default;
+
+  image(image &&rhs) = default;
+
+  image &operator=(const image &rhs) = default;
+
+  image &operator=(image &&rhs) = default;
+
+  ~image() = default;
+
+  template <typename dataT, access::mode accessmode>
+  accessor<dataT, dimensions, accessmode,
+           access::target::image, access::placeholder::false_t>
+  get_access(handler &commandGroupHandler) {
+    return accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t>{};
+  }
+
+  template <typename dataT, access::mode accessmode>
+  accessor<dataT, dimensions, accessmode,
+           access::target::host_image, access::placeholder::false_t>
+  get_access() {
+    return accessor<dataT, dimensions, accessmode, access::target::host_image, access::placeholder::false_t>{};
+  }
+};
+
 } // namespace sycl
 } // namespace cl
diff --git a/clang/test/CodeGenSYCL/image_accessor.cpp b/clang/test/CodeGenSYCL/image_accessor.cpp
new file mode 100644
index 0000000000000..7dca3d7ea6acc
--- /dev/null
+++ b/clang/test/CodeGenSYCL/image_accessor.cpp
@@ -0,0 +1,111 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -I %S/Inputs  -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o %t.ll
+// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-1DRO
+// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-2DRO
+// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DRO
+// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-1DWO
+// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-2DWO
+// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO
+//
+// CHECK-1DRO: %opencl.image1d_ro_t = type opaque
+// CHECK-1DRO: define spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
+//
+// CHECK-2DRO: %opencl.image2d_ro_t = type opaque
+// CHECK-2DRO: define spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
+//
+// CHECK-3DRO: %opencl.image3d_ro_t = type opaque
+// CHECK-3DRO: define spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
+//
+// CHECK-1DWO: %opencl.image1d_wo_t = type opaque
+// CHECK-1DWO: define spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
+//
+// CHECK-2DWO: %opencl.image2d_wo_t = type opaque
+// CHECK-2DWO: define spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
+//
+// CHECK-3DWO: %opencl.image3d_wo_t = type opaque
+// CHECK-3DWO: define spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
+//
+// TODO: Add tests for the image_array opencl datatype support.
+#include "sycl.hpp"
+
+int main() {
+
+  {
+    cl::sycl::image<1> MyImage1d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<1>(3));
+    cl::sycl::queue Q;
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = MyImage1d.get_access<int, cl::sycl::access::mode::read>(cgh);
+
+      cgh.single_task<class image_accessor1dro>([=]() {
+        Acc.use();
+      });
+    });
+  }
+
+  {
+    cl::sycl::image<2> MyImage2d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<2>(3, 2));
+    cl::sycl::queue Q;
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = MyImage2d.get_access<int, cl::sycl::access::mode::read>(cgh);
+
+      cgh.single_task<class image_accessor2dro>([=]() {
+        Acc.use();
+      });
+    });
+  }
+
+  {
+    cl::sycl::image<3> MyImage3d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<3>(3, 2, 4));
+    cl::sycl::queue Q;
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = MyImage3d.get_access<int, cl::sycl::access::mode::read>(cgh);
+
+      cgh.single_task<class image_accessor3dro>([=]() {
+        Acc.use();
+      });
+    });
+  }
+
+  {
+    cl::sycl::image<1> MyImage1d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<1>(3));
+    cl::sycl::queue Q;
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = MyImage1d.get_access<int, cl::sycl::access::mode::write>(cgh);
+
+      cgh.single_task<class image_accessor1dwo>([=]() {
+        Acc.use();
+      });
+    });
+  }
+
+  {
+    cl::sycl::image<2> MyImage2d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<2>(3, 2));
+    cl::sycl::queue Q;
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = MyImage2d.get_access<int, cl::sycl::access::mode::write>(cgh);
+
+      cgh.single_task<class image_accessor2dwo>([=]() {
+        Acc.use();
+      });
+    });
+  }
+
+  {
+    cl::sycl::image<3> MyImage3d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<3>(3, 2, 4));
+    cl::sycl::queue Q;
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = MyImage3d.get_access<int, cl::sycl::access::mode::write>(cgh);
+
+      cgh.single_task<class image_accessor3dwo>([=]() {
+        Acc.use();
+      });
+    });
+  }
+
+  return 0;
+}
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 6313ab843ea9e..3b2046756c30e 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -51,9 +51,9 @@ struct id {
 
 template <int dim>
 struct _ImplT {
-    range<dim> AccessRange;
-    range<dim> MemRange;
-    id<dim> Offset;
+  range<dim> AccessRange;
+  range<dim> MemRange;
+  id<dim> Offset;
 };
 
 template <typename dataT, access::target accessTarget>
@@ -81,7 +81,7 @@ class accessor {
 
 public:
   void use(void) const {}
-  void use(void*) const {}
+  void use(void *) const {}
   _ImplT<dimensions> impl;
 
 private:
@@ -90,6 +90,65 @@ class accessor {
               range<dimensions> MemRange, id<dimensions> Offset) {}
 };
 
+template <int dimensions, access::mode accessmode, access::target accesstarget>
+struct opencl_image_type;
+
+#define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \
+  template <>                                                       \
+  struct opencl_image_type<dim, access::mode::accessmode,           \
+                           access::target::Target> {                \
+    using type = __ocl_image##dim##d_##ifarray_##amsuffix##_t;      \
+  };
+
+#define IMAGETY_READ_3_DIM_IMAGE       \
+  IMAGETY_DEFINE(1, read, ro, image, ) \
+  IMAGETY_DEFINE(2, read, ro, image, ) \
+  IMAGETY_DEFINE(3, read, ro, image, )
+
+#define IMAGETY_WRITE_3_DIM_IMAGE       \
+  IMAGETY_DEFINE(1, write, wo, image, ) \
+  IMAGETY_DEFINE(2, write, wo, image, ) \
+  IMAGETY_DEFINE(3, write, wo, image, )
+
+#define IMAGETY_READ_2_DIM_IARRAY                  \
+  IMAGETY_DEFINE(1, read, ro, image_array, array_) \
+  IMAGETY_DEFINE(2, read, ro, image_array, array_)
+
+#define IMAGETY_WRITE_2_DIM_IARRAY                  \
+  IMAGETY_DEFINE(1, write, wo, image_array, array_) \
+  IMAGETY_DEFINE(2, write, wo, image_array, array_)
+
+IMAGETY_READ_3_DIM_IMAGE
+IMAGETY_WRITE_3_DIM_IMAGE
+
+IMAGETY_READ_2_DIM_IARRAY
+IMAGETY_WRITE_2_DIM_IARRAY
+
+template <int dim, access::mode accessmode, access::target accesstarget>
+struct _ImageImplT {
+#ifdef __SYCL_DEVICE_ONLY__
+  typename opencl_image_type<dim, accessmode, accesstarget>::type MImageObj;
+#else
+  range<dim> AccessRange;
+  range<dim> MemRange;
+  id<dim> Offset;
+#endif
+};
+
+template <typename dataT, int dimensions, access::mode accessmode>
+class accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
+public:
+  void use(void) const {}
+  template <typename... T>
+  void use(T... args) {}
+  template <typename... T>
+  void use(T... args) const {}
+  _ImageImplT<dimensions, accessmode, access::target::image> impl;
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(typename opencl_image_type<dimensions, accessmode, access::target::image>::type ImageObj) { impl.MImageObj = ImageObj; }
+#endif
+};
+
 struct sampler_impl {
 #ifdef __SYCL_DEVICE_ONLY__
   __ocl_sampler_t m_Sampler;
diff --git a/clang/test/SemaSYCL/accessors-targets-image.cpp b/clang/test/SemaSYCL/accessors-targets-image.cpp
new file mode 100644
index 0000000000000..690adeb5e9518
--- /dev/null
+++ b/clang/test/SemaSYCL/accessors-targets-image.cpp
@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct kernel wrapper arguments for
+// image accessors targets.
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+
+  accessor<int, 1, access::mode::read,
+           access::target::image, access::placeholder::false_t>
+      image_acc1d_read;
+  kernel<class use_image1d_r>(
+      [=]() {
+        image_acc1d_read.use();
+      });
+
+  accessor<int, 2, access::mode::read,
+           access::target::image, access::placeholder::false_t>
+      image_acc2d_read;
+  kernel<class use_image2d_r>(
+      [=]() {
+        image_acc2d_read.use();
+      });
+
+  accessor<int, 3, access::mode::read,
+           access::target::image, access::placeholder::false_t>
+      image_acc3d_read;
+  kernel<class use_image3d_r>(
+      [=]() {
+        image_acc3d_read.use();
+      });
+
+  accessor<int, 1, access::mode::write,
+           access::target::image, access::placeholder::false_t>
+      image_acc1d_write;
+  kernel<class use_image1d_w>(
+      [=]() {
+        image_acc1d_write.use();
+      });
+
+  accessor<int, 2, access::mode::write,
+           access::target::image, access::placeholder::false_t>
+      image_acc2d_write;
+  kernel<class use_image2d_w>(
+      [=]() {
+        image_acc2d_write.use();
+      });
+
+  accessor<int, 3, access::mode::write,
+           access::target::image, access::placeholder::false_t>
+      image_acc3d_write;
+  kernel<class use_image3d_w>(
+      [=]() {
+        image_acc3d_write.use();
+      });
+}
+
+// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t)'
+// CHECK: {{.*}}use_image2d_r 'void (__read_only image2d_t)'
+// CHECK: {{.*}}use_image3d_r 'void (__read_only image3d_t)'
+// CHECK: {{.*}}use_image1d_w 'void (__write_only image1d_t)'
+// CHECK: {{.*}}use_image2d_w 'void (__write_only image2d_t)'
+// CHECK: {{.*}}use_image3d_w 'void (__write_only image3d_t)'
diff --git a/sycl/include/CL/__spirv/spirv_types.hpp b/sycl/include/CL/__spirv/spirv_types.hpp
index 6ab6c08ef4a26..e6dfc0eab556c 100644
--- a/sycl/include/CL/__spirv/spirv_types.hpp
+++ b/sycl/include/CL/__spirv/spirv_types.hpp
@@ -57,4 +57,16 @@ enum class GroupOperation : uint32_t {
 #ifndef __SYCL_DEVICE_ONLY__
 typedef void* __ocl_event_t;
 typedef void* __ocl_sampler_t;
+// Adding only the datatypes that can be currently used in SYCL,
+// as per SYCL spec 1.2.1
+typedef void *__ocl_image1d_ro_t;
+typedef void *__ocl_image2d_ro_t;
+typedef void *__ocl_image3d_ro_t;
+typedef void *__ocl_image1d_wo_t;
+typedef void *__ocl_image2d_wo_t;
+typedef void *__ocl_image3d_wo_t;
+typedef void *__ocl_image1d_array_ro_t;
+typedef void *__ocl_image2d_array_ro_t;
+typedef void *__ocl_image1d_array_wo_t;
+typedef void *__ocl_image2d_array_wo_t;
 #endif
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 454e13a00decd..e7d3b9e166130 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -8,10 +8,12 @@
 
 #pragma once
 
+#include <CL/__spirv/spirv_types.hpp>
 #include <CL/sycl/atomic.hpp>
 #include <CL/sycl/buffer.hpp>
 #include <CL/sycl/detail/accessor_impl.hpp>
 #include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/image_ocl_types.hpp>
 #include <CL/sycl/handler.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/pointers.hpp>
diff --git a/sycl/include/CL/sycl/detail/image_ocl_types.hpp b/sycl/include/CL/sycl/detail/image_ocl_types.hpp
new file mode 100644
index 0000000000000..77daa70ccb821
--- /dev/null
+++ b/sycl/include/CL/sycl/detail/image_ocl_types.hpp
@@ -0,0 +1,66 @@
+//===-- Image_ocl_types.hpp - Image OpenCL types --------- ------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// This file is to declare the structs with type as appropriate opencl image
+// types based on Dims, AccessMode and AccessTarget.
+// The macros essentially expand to -
+// template <>
+// struct opencl_image_type<1, access::mode::read, access::target::image> {
+//   using type = __ocl_image1d_ro_t;
+// };
+// 
+// template <>
+// struct opencl_image_type<1, access::mode::write, access::target::image> {
+//   using type = __ocl_image1d_array_wo_t;
+// };
+// 
+// As an example, this can be
+// used as below:
+// detail::opencl_image_type<2, access::mode::read, access::target::image>::type
+//    MyImage;
+//
+
+namespace cl {
+namespace sycl {
+namespace detail {
+template <int Dimensions, access::mode AccessMode, access::target AccessTarget>
+struct opencl_image_type;
+
+#define IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, Ifarray_)            \
+  template <>                                                                  \
+  struct opencl_image_type<Dim, access::mode::AccessMode,                      \
+                           access::target::Target> {                           \
+    using type = __ocl_image##Dim##d_##Ifarray_##AMSuffix##_t;                 \
+  };
+
+#define IMAGETY_READ_3_DIM_IMAGE                                               \
+  IMAGETY_DEFINE(1, read, ro, image, )                                         \
+  IMAGETY_DEFINE(2, read, ro, image, )                                         \
+  IMAGETY_DEFINE(3, read, ro, image, )
+
+#define IMAGETY_WRITE_3_DIM_IMAGE                                              \
+  IMAGETY_DEFINE(1, write, wo, image, )                                        \
+  IMAGETY_DEFINE(2, write, wo, image, )                                        \
+  IMAGETY_DEFINE(3, write, wo, image, )
+
+#define IMAGETY_READ_2_DIM_IARRAY                                              \
+  IMAGETY_DEFINE(1, read, ro, image_array, array_)                             \
+  IMAGETY_DEFINE(2, read, ro, image_array, array_)
+
+#define IMAGETY_WRITE_2_DIM_IARRAY                                             \
+  IMAGETY_DEFINE(1, write, wo, image_array, array_)                            \
+  IMAGETY_DEFINE(2, write, wo, image_array, array_)
+
+IMAGETY_READ_3_DIM_IMAGE
+IMAGETY_WRITE_3_DIM_IMAGE
+
+IMAGETY_READ_2_DIM_IARRAY
+IMAGETY_WRITE_2_DIM_IARRAY
+
+} // namespace detail
+} // namespace sycl
+} // namespace cl