[SYCL] Enable Open CL types required for implementing the SYCL headers.

[erichkeane]

Summary

This patch implements a few of the OpenCL types for SYCL, however
doesn't bother handling semantic analysis as these are not intended to
be used anywhere but in the SYCL implementation.

Signed-off-by: Erich Keane <[e-mail]>

[erichkeane]

@bader : Is this what you meant? Feel free to add any reviewers you want for this. Note it is still dependent on D71016.

[bader]

In D77220#1954909, @erichkeane wrote:

@bader : Is this what you meant? Feel free to add any reviewers you want for this. Note it is still dependent on D71016.

Thanks! I'll add more reviewers.

[keryell]

clang/lib/Sema/SemaSYCL.cpp

357

Curious. What is the purpose?

426

Can it be const?
Can the data structure be partially shared with the one from the previous function?

[erichkeane]

clang/lib/Sema/SemaSYCL.cpp

357

Ah, this is a fix that should actually be in the parent commit: https://reviews.llvm.org/D71016

Since that one was put up for review, the implicit conversion here was removed. I likely forgot to 'git' this out of this and into the cherry pick of the parent revision in my repo. I'll 'remove' it from this patch.

426

It can absolutely be const (change made in the incoming patch), but I don't see a way to combine the two without making matchQualifiedTypeName awkward.

[Anastasia]

Seems reasonable from the OpenCL side.

clang/lib/AST/ASTContext.cpp

void ASTContext::InitBuiltinTypes(const TargetInfo &Target,
   LongDoubleComplexTy = getComplexType(LongDoubleTy);
   Float128ComplexTy   = getComplexType(Float128Ty);
 
   // Builtin types for 'id', 'Class', and 'SEL'.
   InitBuiltinType(ObjCBuiltinIdTy, BuiltinType::ObjCId);
   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);
 #include "clang/Basic/OpenCLImageTypes.def"
 
     InitBuiltinType(OCLSamplerTy, BuiltinType::OCLSampler);
     InitBuiltinType(OCLEventTy, BuiltinType::OCLEvent);
     InitBuiltinType(OCLClkEventTy, BuiltinType::OCLClkEvent);

clang/lib/CodeGen/CodeGenModule.cpp

CodeGenModule::CodeGenModule(ASTContext &C, const HeaderSearchOptions &HSO,
   AllocaInt8PtrTy = Int8Ty->getPointerTo(
       M.getDataLayout().getAllocaAddrSpace());
   ASTAllocaAddressSpace = getTargetCodeGenInfo().getASTAllocaAddressSpace();
 
   RuntimeCC = getTargetCodeGenInfo().getABIInfo().getRuntimeCC();
 
   if (LangOpts.ObjC)
     createObjCRuntime();
-  if (LangOpts.OpenCL)
+  if (LangOpts.OpenCL || LangOpts.SYCLIsDevice)
     createOpenCLRuntime();
   if (LangOpts.OpenMP)
     createOpenMPRuntime();
   if (LangOpts.CUDA)
     createCUDARuntime();
 
   // Enable TBAA unless it's suppressed. ThreadSanitizer needs TBAA even at O0.
   if (LangOpts.Sanitize.has(SanitizerKind::Thread) ||

clang/lib/Sema/Sema.cpp

if (getLangOpts().MSVCCompat) {
     if (getLangOpts().CPlusPlus &&
         IdResolver.begin(&Context.Idents.get("type_info")) == IdResolver.end())
       PushOnScopeChains(Context.buildImplicitRecord("type_info", TTK_Class),
                         TUScope);
 
     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) {
     getOpenCLOptions().addSupport(
         Context.getTargetInfo().getSupportedOpenCLOpts());
     getOpenCLOptions().enableSupportedCore(getLangOpts());
     addImplicitTypedef("sampler_t", Context.OCLSamplerTy);
     addImplicitTypedef("event_t", Context.OCLEventTy);

clang/lib/Sema/SemaInit.cpp

static bool tryObjCWritebackConversion(Sema &S,
   Sequence.AddPassByIndirectCopyRestoreStep(Entity.getType(), ShouldCopy);
   return true;
 }
 
 static bool TryOCLSamplerInitialization(Sema &S,
                                         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()))
     return false;
 
   Sequence.AddOCLSamplerInitStep(DestType);
   return true;
 }
 
 static bool IsZeroInitializer(Expr *Initializer, Sema &S) {
   return Initializer->isIntegerConstantExpr(S.getASTContext()) &&

clang/lib/Sema/SemaSYCL.cpp

 /// Various utilities.
 class Util {
 public:
   using DeclContextDesc = std::pair<clang::Decl::Kind, StringRef>;
 
   /// 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.
   /// \param Ty         the clang type being checked
   /// \param Scopes     the declaration scopes leading from the type to the
   ///     translation unit (excluding the latter)
   static bool matchQualifiedTypeName(const QualType &Ty,
                                      ArrayRef<Util::DeclContextDesc> Scopes);
for (auto Field : LC->fields()) {
       //   - Kernel object field has a scalar type. In this case we should add
       //     simple initialization.
       //   - Kernel object field has a structure or class type. Same handling as
       //     a scalar.
       QualType FieldType = Field->getType();
       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 =
             InitializationKind::CreateDefault(SourceLocation());
         InitializationSequence InitSeq(S, Entity, InitKind, None);
         ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, None);
         InitExprs.push_back(MemberInit.get());
         getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef);
static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
   //     __init method.
   //   - Kernel object field has a scalar type. In this case we should add
   //     kernel parameter with the same type.
   //   - Kernel object field has a structure or class type. Same handling as a
   //     scalar but we should check if this structure/class contains accessors
   //     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);
     else if (ArgTy->isScalarType())
       CreateAndAddPrmDsc(Fld, ArgTy);
     else
       llvm_unreachable("Unsupported kernel parameter type");
   }
 }
 
 // Creates a mangled kernel name for given kernel name type
 static std::string constructKernelName(QualType KernelNameType,
                                        MangleContext &MC) {
   SmallString<256> Result;
   llvm::raw_svector_ostream Out(Result);
 
   MC.mangleTypeName(KernelNameType, Out);
   return std::string(Out.str());
 }
 
 // Generates the OpenCL kernel using KernelCallerFunc (kernel caller
 // function) defined is SYCL headers.
 // Generated OpenCL kernel contains the body of the kernel caller function,
 // receives OpenCL like parameters and additionally does some manipulation to
 // initialize captured lambda/functor fields with these parameters.
 // SYCL runtime marks kernel caller function with sycl_kernel attribute.
bool Util::isSyclAccessorType(const QualType &Ty) {
   static std::array<DeclContextDesc, 3> Scopes = {
       Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
       Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
       Util::DeclContextDesc{clang::Decl::Kind::ClassTemplateSpecialization,
                             "accessor"}};
   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
   // itself. At each step check the context is of expected kind
   // (namespace) and name.
   const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
 
   if (!RecTy)

clang/test/CodeGenSYCL/Inputs/sycl.hpp

 #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 {
   global_buffer = 2014,
   constant_buffer,
   local,
   image,
   host_buffer,

 enum class address_space : int {
   private_space = 0,
   global_space,
   constant_space,
   local_space
 };
 } // 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>
   id(T... args) {} // fake constructor
 private:
   // Some fake field added to see using of id arguments in the
   // kernel wrapper
   int Data;
 };
 
 template <int dim>
 struct range {
   template <typename... T>
   range(T... args) {} // fake constructor
 private:
   // Some fake field added to see using of range arguments in the
   // kernel wrapper
   int Data;
 };
 
 template <int dim>
+struct nd_range {
+};
+
+template <int dim>
 struct _ImplT {
   range<dim> AccessRange;
   range<dim> MemRange;
   id<dim> Offset;
 };
 
 template <typename dataT, int dimensions, access::mode accessmode,
           access::target accessTarget = access::target::global_buffer,
           access::placeholder isPlaceholder = access::placeholder::false_t>
 class accessor {
 
 public:
   void use(void) const {}
   template <typename... T>
   void use(T... args) {}
   template <typename... T>
   void use(T... args) const {}
   _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

clang/test/CodeGenSYCL/image_accessor.cpp

+// 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;
+}

clang/test/CodeGenSYCL/sampler.cpp

+// 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;
+}

clang/test/SemaSYCL/Inputs/sycl.hpp

-#pragma once
+#ifndef SYCL_HPP
+#define SYCL_HPP
+
+// Shared code for SYCL tests
 
 namespace cl {
 namespace sycl {
 namespace access {
 
 enum class target {
   global_buffer = 2014,
   constant_buffer,
public:
   _ImplT<dimensions> impl;
 
 private:
   using PtrType = typename DeviceValueType<dataT, accessTarget>::type *;
   void __init(PtrType 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
+};
+
+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

clang/test/SemaSYCL/accessors-targets-image.cpp

+// 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

clang/test/SemaSYCL/sampler.cpp

+// 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'