[SYCL] Enabling image_array.

The patch maps N dimensional image arrays accessors to N + 1 image
accessors.

operator[](size_t index) method of an N dimensional image array accessor
should return an instance of __image_array_slice__ which provides
interface of N dimensional image accessor to image specified by "index".

The patch introduces this __image_array_slice__ class which holds
original accessor and adjusts read, write, get_size and get_count
methods to imitate operations with N dimensional image.

This mapping approach is an easiest way to implement image array accesors

Signed-off-by: Vlad Romanov <[email protected]>

Author: romanovvlad

sycl/include/CL/sycl/accessor.hpp

@@ -11,6 +11,7 @@
 #include <CL/__spirv/spirv_types.hpp>
 #include <CL/sycl/atomic.hpp>
 #include <CL/sycl/buffer.hpp>
+#include <CL/sycl/exception.hpp>
 #include <CL/sycl/detail/accessor_impl.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/generic_type_traits.hpp>
@@ -242,6 +243,26 @@ class accessor_common {
   };
 };
 
+template <int Dim, typename T> struct IsValidCoordDataT;
+template <typename T> struct IsValidCoordDataT<1, T> {
+  constexpr static bool value =
+      detail::is_contained<T, detail::type_list<cl_int, cl_float>>::type::value;
+};
+template <typename T> struct IsValidCoordDataT<2, T> {
+  constexpr static bool value =
+      detail::is_contained<T,
+                           detail::type_list<cl_int2, cl_float2>>::type::value;
+};
+template <typename T> struct IsValidCoordDataT<3, T> {
+  constexpr static bool value =
+      detail::is_contained<T,
+                           detail::type_list<cl_int4, cl_float4>>::type::value;
+};
+
+template <typename DataT, int Dimensions, access::mode AccessMode,
+          access::placeholder IsPlaceholder>
+class __image_array_slice__;
+
 // Image accessor
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::target AccessTarget, access::placeholder IsPlaceholder>
@@ -266,6 +287,9 @@ class image_accessor
 
 private:
 #endif
+  template <typename T1, int T2, access::mode T3, access::placeholder T4>
+  friend class __image_array_slice__;
+
   constexpr static bool IsHostImageAcc =
       (AccessTarget == access::target::host_image);
 
@@ -310,6 +334,30 @@ class image_accessor
       throw feature_not_supported("Images are not supported by this device.");
   }
 
+#ifdef __SYCL_DEVICE_ONLY__
+
+  sycl::vec<int, Dimensions> getCountInternal() const {
+    return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(MImageObj);
+  }
+
+  size_t getElementSize() const {
+    int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
+    int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
+    int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
+    return ElementSize;
+  }
+
+#else
+
+  sycl::vec<int, Dimensions> getCountInternal() const {
+    // TODO: Implement for host.
+    throw runtime_error(
+        "image::getCountInternal() is not implemented for host");
+    return sycl::vec<int, Dimensions>{1};
+  }
+
+#endif
+
 public:
   using value_type = DataT;
   using reference = DataT &;
@@ -371,30 +419,6 @@ class image_accessor
   }
 #endif
 
-  template <typename AllocatorT, int Dims = Dimensions,
-            typename = detail::enable_if_t<(Dims > 0) && (Dims < 3) &&
-                                           IsImageArrayAcc>>
-  image_accessor(image<Dims + 1, AllocatorT> &ImageRef,
-                 handler &CommandGroupHandlerRef, int ImageElementSize)
-#ifdef __SYCL_DEVICE_ONLY__
-  {
-    // No implementation needed for device. The constructor is only called by
-    // host.
-  }
-#else
-      : AccessorBaseHost(id<3>(0, 0, 0) /* Offset,*/,
-                         detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
-                         detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
-                         AccessMode, detail::getSyclObjImpl(ImageRef).get(),
-                         Dimensions, ImageElementSize),
-        MImageCount(ImageRef.get_count()),
-        MImageSize(MImageCount * ImageElementSize) {
-    checkDeviceFeatureSupported<info::device::image_support>(
-        CommandGroupHandlerRef.MQueue->get_device());
-    // TODO: Implement this function.
-  }
-#endif
-
   /* -- common interface members -- */
 
   // operator == and != need to be defined only for host application as per the
@@ -426,37 +450,21 @@ class image_accessor
 
   template <int Dims = Dimensions> size_t get_count() const;
 
-  template <> size_t get_count<1>() const {
-    return __invoke_ImageQuerySize<int, OCLImageTy>(MImageObj);
-  }
+  template <> size_t get_count<1>() const { return getCountInternal(); }
   template <> size_t get_count<2>() const {
-    cl_int2 Count = __invoke_ImageQuerySize<cl_int2, OCLImageTy>(MImageObj);
+    cl_int2 Count = getCountInternal();
     return (Count.x() * Count.y());
   };
   template <> size_t get_count<3>() const {
-    cl_int3 Count = __invoke_ImageQuerySize<cl_int3, OCLImageTy>(MImageObj);
+    cl_int3 Count = getCountInternal();
     return (Count.x() * Count.y() * Count.z());
   };
+
 #else
   size_t get_size() const { return MImageSize; };
   size_t get_count() const { return MImageCount; };
 #endif
 
-  template <int Dim, typename T> struct IsValidCoordDataT;
-  template <typename T> struct IsValidCoordDataT<1, T> {
-    constexpr static bool value =
-        detail::is_contained<T,
-                             detail::type_list<cl_int, cl_float>>::type::value;
-  };
-  template <typename T> struct IsValidCoordDataT<2, T> {
-    constexpr static bool value = detail::is_contained<
-        T, detail::type_list<cl_int2, cl_float2>>::type::value;
-  };
-  template <typename T> struct IsValidCoordDataT<3, T> {
-    constexpr static bool value = detail::is_contained<
-        T, detail::type_list<cl_int4, cl_float4>>::type::value;
-  };
-
   // Available only when:
   // (accessTarget == access::target::image && accessMode == access::mode::read)
   // || (accessTarget == access::target::host_image && ( accessMode ==
@@ -514,10 +522,96 @@ class image_accessor
         "Read API is not implemented on host.");
 #endif
   }
+};
+
+template <typename DataT, int Dimensions, access::mode AccessMode,
+          access::placeholder IsPlaceholder>
+class __image_array_slice__ {
+
+  static_assert(Dimensions < 3,
+                "Image slice cannot have more then 2 dimensions");
+
+  constexpr static int AdjustedDims = (Dimensions == 2) ? 4 : Dimensions + 1;
+
+  template <typename CoordT,
+            typename CoordElemType =
+                typename detail::TryToGetElementType<CoordT>::type>
+  sycl::vec<CoordElemType, AdjustedDims>
+  getAdjustedCoords(const CoordT &Coords) const {
+    CoordElemType LastCoord = 0;
+
+    if (std::is_same<float, CoordElemType>::value) {
+      sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getCountInternal();
+      LastCoord =
+          MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
+    } else {
+      LastCoord = MIdx;
+    }
 
-  // Available only when: accessTarget == access::target::image_array &&
-  // dimensions < 3
-  //__image_array_slice__ operator[](size_t index) const;
+    sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
+    sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
+                                                          LeftoverCoords};
+    return AdjustedCoords;
+  }
+
+public:
+  __image_array_slice__(accessor<DataT, Dimensions, AccessMode,
+                           access::target::image_array, IsPlaceholder>
+                      BaseAcc,
+                  size_t Idx)
+      : MBaseAcc(BaseAcc), MIdx(Idx) {}
+
+  template <typename CoordT, int Dims = Dimensions,
+            typename = detail::enable_if_t<
+                (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value)>>
+  DataT read(const CoordT &Coords) const {
+    return MBaseAcc.read(getAdjustedCoords(Coords));
+  }
+
+  template <typename CoordT, int Dims = Dimensions,
+            typename = detail::enable_if_t<
+                (Dims > 0) && IsValidCoordDataT<Dims, CoordT>::value>>
+  DataT read(const CoordT &Coords, const sampler &Smpl) const {
+    return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
+  }
+
+  template <typename CoordT, int Dims = Dimensions,
+            typename = detail::enable_if_t<
+                (Dims > 0) && IsValidCoordDataT<Dims, CoordT>::value>>
+  void write(const CoordT &Coords, const DataT &Color) const {
+    return MBaseAcc.write(getAdjustedCoords(Coords), Color);
+  }
+
+
+#ifdef __SYCL_DEVICE_ONLY__
+  size_t get_size() const { return MBaseAcc.getElementSize() * get_count(); }
+
+  template <int Dims = Dimensions> size_t get_count() const;
+
+  template <> size_t get_count<1>() const {
+    cl_int2 Count = MBaseAcc.getCountInternal();
+    return Count.x();
+  }
+  template <> size_t get_count<2>() const {
+    cl_int3 Count = MBaseAcc.getCountInternal();
+    return (Count.x() * Count.y());
+  };
+#else
+
+  size_t get_size() const {
+    return MBaseAcc.MImageSize / MBaseAcc.getAccessRange()[Dimensions];
+  };
+
+  size_t get_count() const {
+    return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions];
+  };
+#endif
+
+private:
+  size_t MIdx;
+  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
+           IsPlaceholder>
+      MBaseAcc;
 };
 
 } // namespace detail
@@ -1066,6 +1160,7 @@ class accessor<DataT, Dimensions, AccessMode, access::target::host_image,
             Image, (detail::getSyclObjImpl(Image))->getElementSize()) {}
 };
 
+
 // Available only when: accessTarget == access::target::image_array &&
 // dimensions < 3
 // template <typename AllocatorT> accessor(image<dimensions + 1,
@@ -1074,20 +1169,35 @@ template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
 class accessor<DataT, Dimensions, AccessMode, access::target::image_array,
                IsPlaceholder>
-    : public detail::image_accessor<DataT, Dimensions, AccessMode,
-                                    access::target::image_array,
+    : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
+                                    access::target::image,
                                     IsPlaceholder> {
-  // TODO: Add Constructor.
 #ifdef __SYCL_DEVICE_ONLY__
 private:
   using OCLImageTy =
-      typename detail::opencl_image_type<Dimensions, AccessMode,
-                                         access::target::image_array>::type;
+      typename detail::opencl_image_type<Dimensions + 1, AccessMode,
+                                         access::target::image>::type;
 
   // Front End requires this method to be defined in the accessor class.
   // It does not call the base class's init method.
   void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
 #endif
+public:
+  template <typename AllocatorT>
+  accessor(cl::sycl::image<Dimensions + 1, AllocatorT> &Image,
+           handler &CommandGroupHandler)
+      : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
+                               access::target::image, IsPlaceholder>(
+            Image, CommandGroupHandler,
+            (detail::getSyclObjImpl(Image))->getElementSize()) {
+    CommandGroupHandler.associateWithHandler(*this);
+  }
+
+  detail::__image_array_slice__<DataT, Dimensions, AccessMode, IsPlaceholder>
+  operator[](size_t Index) const {
+    return detail::__image_array_slice__<DataT, Dimensions, AccessMode,
+                                   IsPlaceholder>(*this, Index);
+  }
 };
 
 } // namespace sycl

sycl/include/CL/sycl/detail/image_ocl_types.hpp

@@ -136,6 +136,7 @@ inline int getSPIRVNumChannels(int ImageChannelOrder) {
   case 18: // sBGRA
     // TODO: Enable the below assert after assert is supported for device
     // compiler. assert(!"Unhandled image channel order in sycl.");
+  default:
     return 0;
   }
 }