[SYCL] Clang Front End Support for image classes.

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)

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

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

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;