Index: clang/lib/AST/ASTContext.cpp
===================================================================
--- clang/lib/AST/ASTContext.cpp
+++ clang/lib/AST/ASTContext.cpp
@@ -1402,6 +1402,15 @@
   InitBuiltinType(ObjCBuiltinClassTy, BuiltinType::ObjCClass);
   InitBuiltinType(ObjCBuiltinSelTy, BuiltinType::ObjCSel);
 
+  if (LangOpts.SYCLIsDevice) {
+    InitBuiltinType(OCLSamplerTy, BuiltinType::OCLSampler);
+    InitBuiltinType(OCLEventTy, BuiltinType::OCLEvent);
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
+    InitBuiltinType(SingletonId, BuiltinType::Id);
+#include "clang/Basic/OpenCLImageTypes.def"
+#undef IMAGE_TYPE
+  }
+
   if (LangOpts.OpenCL) {
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     InitBuiltinType(SingletonId, BuiltinType::Id);
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -134,7 +134,7 @@
 
   if (LangOpts.ObjC)
     createObjCRuntime();
-  if (LangOpts.OpenCL)
+  if (LangOpts.OpenCL || LangOpts.SYCLIsDevice)
     createOpenCLRuntime();
   if (LangOpts.OpenMP)
     createOpenMPRuntime();
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -288,6 +288,17 @@
     addImplicitTypedef("size_t", Context.getSizeType());
   }
 
+  if (getLangOpts().SYCLIsDevice) {
+    addImplicitTypedef("__ocl_sampler_t", Context.OCLSamplerTy);
+    addImplicitTypedef("__ocl_event_t", Context.OCLEventTy);
+#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)
   // core features.
   if (getLangOpts().OpenCL) {
Index: clang/lib/Sema/SemaInit.cpp
===================================================================
--- clang/lib/Sema/SemaInit.cpp
+++ clang/lib/Sema/SemaInit.cpp
@@ -5463,9 +5463,10 @@
                                         InitializationSequence &Sequence,
                                         QualType DestType,
                                         Expr *Initializer) {
-  if (!S.getLangOpts().OpenCL || !DestType->isSamplerT() ||
+  if ((!S.getLangOpts().OpenCL && !S.getLangOpts().SYCLIsDevice) ||
+      !DestType->isSamplerT() ||
       (!Initializer->isIntegerConstantExpr(S.Context) &&
-      !Initializer->getType()->isSamplerT()))
+       !Initializer->getType()->isSamplerT()))
     return false;
 
   Sequence.AddOCLSamplerInitStep(DestType);
Index: clang/lib/Sema/SemaSYCL.cpp
===================================================================
--- clang/lib/Sema/SemaSYCL.cpp
+++ clang/lib/Sema/SemaSYCL.cpp
@@ -27,6 +27,9 @@
   /// Checks whether given clang type is a full specialization of the SYCL
   /// accessor class.
   static bool isSyclAccessorType(const QualType &Ty);
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// sampler class.
+  static bool isSyclSamplerType(const QualType &Ty);
 
   /// Checks whether given clang type is declared in the given hierarchy of
   /// declaration contexts.
@@ -221,7 +224,8 @@
       CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl();
       InitializedEntity Entity =
           InitializedEntity::InitializeMember(Field, &VarEntity);
-      if (Util::isSyclAccessorType(FieldType)) {
+      if (Util::isSyclAccessorType(FieldType) ||
+          Util::isSyclSamplerType(FieldType)) {
         // Initialize kernel object field with the default constructor and
         // construct a call of __init method.
         InitializationKind InitKind =
@@ -332,7 +336,7 @@
   //     and add parameter decriptor for them properly.
   for (const auto *Fld : KernelObj->fields()) {
     QualType ArgTy = Fld->getType();
-    if (Util::isSyclAccessorType(ArgTy))
+    if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy))
       createSpecialSYCLObjParamDesc(Fld, ArgTy);
     else if (ArgTy->isStructureOrClassType())
       CreateAndAddPrmDsc(Fld, ArgTy);
@@ -418,6 +422,15 @@
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclSamplerType(const QualType &Ty) {
+  static const std::array<DeclContextDesc, 3> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::CXXRecord,
+                            "sampler"}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::matchQualifiedTypeName(const QualType &Ty,
                                   ArrayRef<Util::DeclContextDesc> Scopes) {
   // The idea: check the declaration context chain starting from the type
Index: clang/test/CodeGenSYCL/Inputs/sycl.hpp
===================================================================
--- clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -1,7 +1,32 @@
 #pragma once
 
+#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
+
+// Dummy runtime classes to model SYCL API.
 namespace cl {
 namespace sycl {
+struct sampler_impl {
+#ifdef __SYCL_DEVICE_ONLY__
+  __ocl_sampler_t m_Sampler;
+#endif
+};
+
+class sampler {
+  struct sampler_impl impl;
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
+#endif
+
+public:
+  void use(void) const {}
+};
+
+template <int dimensions = 1>
+class group {
+public:
+  group() = default; // fake constructor
+};
+
 namespace access {
 
 enum class target {
@@ -36,6 +61,39 @@
 };
 } // namespace access
 
+namespace property {
+
+enum prop_type {
+  use_host_ptr = 0,
+  use_mutex,
+  context_bound,
+  enable_profiling,
+  base_prop
+};
+
+struct property_base {
+  virtual prop_type type() const = 0;
+};
+} // namespace property
+
+class property_list {
+public:
+  template <typename... propertyTN>
+  property_list(propertyTN... props) {}
+
+  template <typename propertyT>
+  bool has_property() const { return true; }
+
+  template <typename propertyT>
+  propertyT get_property() const {
+    return propertyT{};
+  }
+
+  bool operator==(const property_list &rhs) const { return false; }
+
+  bool operator!=(const property_list &rhs) const { return false; }
+};
+
 template <int dim>
 struct id {
   template <typename... T>
@@ -57,6 +115,10 @@
 };
 
 template <int dim>
+struct nd_range {
+};
+
+template <int dim>
 struct _ImplT {
   range<dim> AccessRange;
   range<dim> MemRange;
@@ -77,10 +139,294 @@
   _ImplT<dimensions> impl;
 
 private:
-  void __init(__attribute__((opencl_global)) dataT *Ptr,
-              range<dimensions> AccessRange,
+  void __init(__attribute__((opencl_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 {};
+class event {};
+
+class queue {
+public:
+  template <typename T>
+  event submit(T cgf) { return event{}; }
+
+  void wait() {}
+  void wait_and_throw() {}
+  void throw_asynchronous() {}
+};
+
+class auto_name {};
+template <typename Name, typename Type>
+struct get_kernel_name_t {
+  using name = Name;
+};
+template <typename Type>
+struct get_kernel_name_t<auto_name, Type> {
+  using name = Type;
+};
+#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
+template <typename KernelName = auto_name, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) {
+  kernelFunc();
+}
+
+template <typename KernelName, typename KernelType, int Dims>
+ATTR_SYCL_KERNEL void
+kernel_parallel_for(KernelType KernelFunc) {
+  KernelFunc(id<Dims>());
+}
+
+template <typename KernelName, typename KernelType, int Dims>
+ATTR_SYCL_KERNEL void
+kernel_parallel_for_work_group(KernelType KernelFunc) {
+  KernelFunc(group<Dims>());
+}
+
+class handler {
+public:
+  template <typename KernelName = auto_name, typename KernelType, int Dims>
+  void parallel_for(range<Dims> numWorkItems, KernelType kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for<NameT, KernelType, Dims>(kernelFunc);
+#else
+    kernelFunc();
+#endif
+  }
+
+  template <typename KernelName = auto_name, typename KernelType, int Dims>
+  void parallel_for_work_group(range<Dims> numWorkGroups, range<Dims> WorkGroupSize, KernelType kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_work_group<NameT, KernelType, Dims>(kernelFunc);
+#else
+    group<Dims> G;
+    kernelFunc(G);
+#endif
+  }
+
+  template <typename KernelName = auto_name, typename KernelType>
+  void single_task(KernelType kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_single_task<NameT>(kernelFunc);
+#else
+    kernelFunc();
+#endif
+  }
+};
+
+class stream {
+public:
+  stream(unsigned long BufferSize, unsigned long MaxStatementSize,
+         handler &CGH) {}
+
+  void __init() {}
+
+  void __finalize() {}
+};
+
+template <typename T>
+const stream& operator<<(const stream &S, T&&) {
+  return S;
+}
+
+template <typename T, int dimensions = 1,
+          typename AllocatorT = int /*fake type as AllocatorT is not used*/>
+class buffer {
+public:
+  using value_type = T;
+  using reference = value_type &;
+  using const_reference = const value_type &;
+  using allocator_type = AllocatorT;
+
+  template <typename... ParamTypes>
+  buffer(ParamTypes... args) {} // fake constructor
+
+  buffer(const range<dimensions> &bufferRange,
+         const property_list &propList = {}) {}
+
+  buffer(T *hostData, const range<dimensions> &bufferRange,
+         const property_list &propList = {}) {}
+
+  buffer(const T *hostData, const range<dimensions> &bufferRange,
+         const property_list &propList = {}) {}
+
+  buffer(const buffer &rhs) = default;
+
+  buffer(buffer &&rhs) = default;
+
+  buffer &operator=(const buffer &rhs) = default;
+
+  buffer &operator=(buffer &&rhs) = default;
+
+  ~buffer() = default;
+
+  range<dimensions> get_range() const { return range<dimensions>{}; }
+
+  template <access::mode mode,
+            access::target target = access::target::global_buffer>
+  accessor<T, dimensions, mode, target, access::placeholder::false_t>
+  get_access(handler &commandGroupHandler) {
+    return accessor<T, dimensions, mode, target, access::placeholder::false_t>{};
+  }
+
+  template <access::mode mode>
+  accessor<T, dimensions, mode, access::target::host_buffer,
+           access::placeholder::false_t>
+  get_access() {
+    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
Index: clang/test/CodeGenSYCL/image_accessor.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenSYCL/image_accessor.cpp
@@ -0,0 +1,111 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -I %S/Inputs -fsycl -fsycl-is-device -disable-llvm-passes -emit-llvm %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;
+}
Index: clang/test/CodeGenSYCL/sampler.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenSYCL/sampler.cpp
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -I %S/Inputs -fsycl -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck --enable-var-scope %s
+// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
+// CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
+// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8
+// CHECK-NEXT: [[GEP:%[0-9]+]]  = getelementptr inbounds %class.anon, %class.anon* [[ANON]], i32 0, i32 0
+// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8
+// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler"* [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
+
+#include "sycl.hpp"
+
+template <typename KernelName, typename KernelType>
+__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  cl::sycl::sampler smplr;
+  kernel_single_task<class first_kernel>([=]() {
+    smplr.use();
+  });
+
+  return 0;
+}
Index: clang/test/SemaSYCL/Inputs/sycl.hpp
===================================================================
--- clang/test/SemaSYCL/Inputs/sycl.hpp
+++ clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -1,4 +1,7 @@
-#pragma once
+#ifndef SYCL_HPP
+#define SYCL_HPP
+
+// Shared code for SYCL tests
 
 namespace cl {
 namespace sycl {
@@ -83,5 +86,115 @@
               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;
+#endif
+};
+
+class sampler {
+  struct sampler_impl impl;
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
+#endif
+
+public:
+  void use(void) const {}
+};
+
+class event {};
+class queue {
+public:
+  template <typename T>
+  event submit(T cgf) { return event{}; }
+};
+class auto_name {};
+template <typename Name, typename Type>
+struct get_kernel_name_t {
+  using name = Name;
+};
+template <typename Type>
+struct get_kernel_name_t<auto_name, Type> {
+  using name = Type;
+};
+#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
+template <typename KernelName = auto_name, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) {
+  kernelFunc();
+}
+class handler {
+public:
+  template <typename KernelName = auto_name, typename KernelType>
+  void single_task(KernelType kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_single_task<NameT>(kernelFunc);
+#else
+    kernelFunc();
+#endif
+  }
+};
+
 } // namespace sycl
 } // namespace cl
+
+#endif
Index: clang/test/SemaSYCL/accessors-targets-image.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/accessors-targets-image.cpp
@@ -0,0 +1,74 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -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)'
+
+// TODO: SYCL specific fail - analyze and enable
+// XFAIL: windows-msvc
Index: clang/test/SemaSYCL/sampler.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/sampler.cpp
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+
+#include <sycl.hpp>
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  cl::sycl::sampler Sampler;
+  kernel<class use_kernel_for_test>([=]() {
+    Sampler.use();
+  });
+  return 0;
+}
+
+// Check declaration of the test kernel
+// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (sampler_t)'
+//
+// Check parameters of the test kernel
+// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] 'sampler_t'
+//
+// Check that sampler field of the test kernel object is initialized using __init method
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__ocl_sampler_t)' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::sampler':'cl::sycl::sampler' lvalue
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})'
+//
+// Check the parameters of __init method
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' 'sampler_t'