[SYCL] Add host_task image accessor support (#9718)

This commit adds the support for SYCL 2020 image accessors in host_task
commands as well as implementations of their member functions on host.

---------

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/accessor.hpp

@@ -19,15 +19,22 @@ class handler;
 namespace detail {
 
 class AccessorBaseHost;
+class UnsampledImageAccessorBaseHost;
+class SampledImageAccessorBaseHost;
 
 #ifdef __SYCL_DEVICE_ONLY__
-// In device compilation accessor isn't inherited from AccessorBaseHost, so
+// In device compilation accessor isn't inherited from host base classes, so
 // can't detect by it. Since we don't expect it to be ever called in device
 // execution, just use blind void *.
 inline void associateWithHandler(handler &, void *, access::target) {}
+inline void associateWithHandler(handler &, void *, image_target) {}
 #else
 __SYCL_EXPORT void associateWithHandler(handler &, AccessorBaseHost *,
                                         access::target);
+__SYCL_EXPORT void
+associateWithHandler(handler &, UnsampledImageAccessorBaseHost *, image_target);
+__SYCL_EXPORT void
+associateWithHandler(handler &, SampledImageAccessorBaseHost *, image_target);
 #endif
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)

sycl/include/sycl/detail/handler_proxy.hpp

@@ -102,11 +102,23 @@ template <typename DataT, int Dimensions>
 struct owner_less<local_accessor<DataT, Dimensions>>
     : public detail::owner_less_base<local_accessor<DataT, Dimensions>> {};
 
+template <typename DataT, int Dimensions, access_mode AccessMode,
+          image_target AccessTarget>
+struct owner_less<
+    unsampled_image_accessor<DataT, Dimensions, AccessMode, AccessTarget>>
+    : public detail::owner_less_base<unsampled_image_accessor<
+          DataT, Dimensions, AccessMode, AccessTarget>> {};
+
 template <typename DataT, int Dimensions, access_mode AccessMode>
 struct owner_less<host_unsampled_image_accessor<DataT, Dimensions, AccessMode>>
     : public detail::owner_less_base<
           host_unsampled_image_accessor<DataT, Dimensions, AccessMode>> {};
 
+template <typename DataT, int Dimensions, image_target AccessTarget>
+struct owner_less<sampled_image_accessor<DataT, Dimensions, AccessTarget>>
+    : public detail::owner_less_base<
+          sampled_image_accessor<DataT, Dimensions, AccessTarget>> {};
+
 template <typename DataT, int Dimensions>
 struct owner_less<host_sampled_image_accessor<DataT, Dimensions>>
     : public detail::owner_less_base<

sycl/include/sycl/ext/oneapi/owner_less.hpp

@@ -454,13 +454,20 @@ class __SYCL_EXPORT handler {
   bool is_host() { return MIsHost; }
 
 #ifdef __SYCL_DEVICE_ONLY__
-  // In device compilation accessor isn't inherited from AccessorBaseHost, so
+  // In device compilation accessor isn't inherited from host base classes, so
   // can't detect by it. Since we don't expect it to be ever called in device
   // execution, just use blind void *.
   void associateWithHandler(void *AccBase, access::target AccTarget);
+  void associateWithHandler(void *AccBase, image_target AccTarget);
 #else
+  void associateWithHandlerCommon(detail::AccessorImplPtr AccImpl,
+                                  int AccTarget);
   void associateWithHandler(detail::AccessorBaseHost *AccBase,
                             access::target AccTarget);
+  void associateWithHandler(detail::UnsampledImageAccessorBaseHost *AccBase,
+                            image_target AccTarget);
+  void associateWithHandler(detail::SampledImageAccessorBaseHost *AccBase,
+                            image_target AccTarget);
 #endif
 
   // Recursively calls itself until arguments pack is fully processed.
@@ -2935,6 +2942,10 @@ class __SYCL_EXPORT handler {
   friend void detail::associateWithHandler(handler &,
                                            detail::AccessorBaseHost *,
                                            access::target);
+  friend void detail::associateWithHandler(
+      handler &, detail::UnsampledImageAccessorBaseHost *, image_target);
+  friend void detail::associateWithHandler(
+      handler &, detail::SampledImageAccessorBaseHost *, image_target);
 #endif
 
   friend class ::MockHandler;

sycl/include/sycl/handler.hpp

@@ -19,6 +19,16 @@ void associateWithHandler(handler &CGH, AccessorBaseHost *Acc,
   CGH.associateWithHandler(Acc, Target);
 }
 
+void associateWithHandler(handler &CGH, UnsampledImageAccessorBaseHost *Acc,
+                          image_target Target) {
+  CGH.associateWithHandler(Acc, Target);
+}
+
+void associateWithHandler(handler &CGH, SampledImageAccessorBaseHost *Acc,
+                          image_target Target) {
+  CGH.associateWithHandler(Acc, Target);
+}
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/source/detail/handler_proxy.cpp

@@ -376,9 +376,8 @@ void handler::addReduction(const std::shared_ptr<const void> &ReduObj) {
   MImpl->MAuxiliaryResources.push_back(ReduObj);
 }
 
-void handler::associateWithHandler(detail::AccessorBaseHost *AccBase,
-                                   access::target AccTarget) {
-  detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase);
+void handler::associateWithHandlerCommon(detail::AccessorImplPtr AccImpl,
+                                         int AccTarget) {
   detail::Requirement *Req = AccImpl.get();
   // Add accessor to the list of requirements.
   CGData.MRequirements.push_back(Req);
@@ -387,8 +386,25 @@ void handler::associateWithHandler(detail::AccessorBaseHost *AccBase,
   // Add an accessor to the handler list of associated accessors.
   // For associated accessors index does not means nothing.
   MAssociatedAccesors.emplace_back(detail::kernel_param_kind_t::kind_accessor,
-                                   Req, static_cast<int>(AccTarget),
-                                   /*index*/ 0);
+                                   Req, AccTarget, /*index*/ 0);
+}
+
+void handler::associateWithHandler(detail::AccessorBaseHost *AccBase,
+                                   access::target AccTarget) {
+  associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase),
+                             static_cast<int>(AccTarget));
+}
+
+void handler::associateWithHandler(
+    detail::UnsampledImageAccessorBaseHost *AccBase, image_target AccTarget) {
+  associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase),
+                             static_cast<int>(AccTarget));
+}
+
+void handler::associateWithHandler(
+    detail::SampledImageAccessorBaseHost *AccBase, image_target AccTarget) {
+  associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase),
+                             static_cast<int>(AccTarget));
 }
 
 static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index,

sycl/source/handler.cpp

@@ -200,7 +200,7 @@ CoordT<ImgT, Dims> DelinearizeToCoord(size_t Idx, range<Dims> ImageRange,
   } else if constexpr (Dims == 2) {
     Out = CoordT<ImgT, Dims>{Idx % ImageRange[0], Idx / ImageRange[0]};
   } else {
-    Out = CoordT<ImgT, Dims>{Idx % ImageRange[0] % ImageRange[1],
+    Out = CoordT<ImgT, Dims>{Idx % ImageRange[0],
                              Idx / ImageRange[0] % ImageRange[1],
                              Idx / ImageRange[0] / ImageRange[1], 0};
   }
@@ -328,3 +328,119 @@ ApplyAddressingMode(CoordT<ImageType::Sampled, Dims> Coord,
     }
   }
 }
+
+template <image_format Format> static constexpr size_t getMaxInt() {
+  using rep_elem_type = typename FormatTraits<Format>::rep_elem_type;
+  return static_cast<size_t>(std::numeric_limits<rep_elem_type>::max());
+}
+
+template <image_format Format>
+typename FormatTraits<Format>::pixel_type PickNewColor(size_t I,
+                                                       size_t AccSize) {
+  using PixelType = typename FormatTraits<Format>::pixel_type;
+  size_t Idx = I * 4;
+
+  // Pick a new color. Make sure it isn't too big for the data type.
+  PixelType NewColor{Idx, Idx + 1, Idx + 2, Idx + 3};
+  NewColor = sycl::min(NewColor, PixelType{getMaxInt<Format>()});
+  if constexpr (FormatTraits<Format>::Normalized)
+    NewColor /= AccSize * 4;
+  return NewColor;
+}
+
+// Implemented as specified by the OpenCL 1.2 specification for
+// CLK_FILTER_NEAREST.
+template <image_format Format, addressing_mode AddrMode, int Dims>
+typename FormatTraits<Format>::pixel_type
+ReadNearest(typename FormatTraits<Format>::rep_elem_type *RefData,
+            CoordT<ImageType::Sampled, Dims> Coord, range<2> ImagePitch,
+            range<Dims> ImageRange, bool Normalized) {
+  CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
+  if constexpr (AddrMode == addressing_mode::repeat) {
+    assert(Normalized);
+    AdjCoord -= sycl::floor(AdjCoord);
+    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
+    AdjCoord = sycl::floor(AdjCoord);
+  } else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
+    assert(Normalized);
+    AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
+    AdjCoord = sycl::fabs(Coord - AdjCoord);
+    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
+    AdjCoord = sycl::floor(AdjCoord);
+  } else {
+    if (Normalized)
+      AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
+    AdjCoord = sycl::floor(AdjCoord);
+  }
+  AdjCoord = ApplyAddressingMode<AddrMode>(AdjCoord, ImageRange);
+  return SimulateRead<Format, ImageType::Sampled>(RefData, AdjCoord, ImagePitch,
+                                                  ImageRange, false);
+}
+
+// Implemented as specified by the OpenCL 1.2 specification for
+// CLK_FILTER_LINEAR.
+template <image_format Format, addressing_mode AddrMode, int Dims>
+float4 CalcLinearRead(typename FormatTraits<Format>::rep_elem_type *RefData,
+                      CoordT<ImageType::Sampled, Dims> Coord,
+                      range<2> ImagePitch, range<Dims> ImageRange,
+                      bool Normalized) {
+  using UpscaledCoordT = CoordT<ImageType::Sampled, 3>;
+
+  auto Read = [&](UpscaledCoordT UpCoord) {
+    auto DownCoord = DownscaleCoord<Dims>(UpCoord);
+    return SimulateRead<Format, ImageType::Sampled>(
+        RefData, DownCoord, ImagePitch, ImageRange, false);
+  };
+
+  CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
+  if constexpr (AddrMode == addressing_mode::repeat) {
+    assert(Normalized);
+    AdjCoord -= floor(AdjCoord);
+    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
+  } else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
+    assert(Normalized);
+    AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
+    AdjCoord = sycl::fabs(Coord - AdjCoord);
+    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
+  } else {
+    if (Normalized)
+      AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
+  }
+
+  auto Prev = sycl::floor(AdjCoord - 0.5f);
+  auto Next = Prev + 1;
+  auto CA = (AdjCoord - 0.5f) - Prev;
+
+  Prev = ApplyAddressingMode<AddrMode>(Prev, ImageRange);
+  Next = ApplyAddressingMode<AddrMode>(Next, ImageRange);
+
+  auto UPrev = UpscaleCoord<Dims>(Prev);
+  auto UNext = UpscaleCoord<Dims>(Next);
+  auto UCA = UpscaleCoord<Dims>(CA, 1);
+
+  auto CA000 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UPrev[2], 0})
+                   .template convert<float>() *
+               (1 - UCA[0]) * (1 - UCA[1]) * (1 - UCA[2]);
+  auto CA100 = Read(UpscaledCoordT{UNext[0], UPrev[1], UPrev[2], 0})
+                   .template convert<float>() *
+               UCA[0] * (1 - UCA[1]) * (1 - UCA[2]);
+  auto CA010 = Read(UpscaledCoordT{UPrev[0], UNext[1], UPrev[2], 0})
+                   .template convert<float>() *
+               (1 - UCA[0]) * UCA[1] * (1 - UCA[2]);
+  auto CA110 = Read(UpscaledCoordT{UNext[0], UNext[1], UPrev[2], 0})
+                   .template convert<float>() *
+               UCA[0] * UCA[1] * (1 - UCA[2]);
+  auto CA001 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UNext[2], 0})
+                   .template convert<float>() *
+               (1 - UCA[0]) * (1 - UCA[1]) * UCA[2];
+  auto CA101 = Read(UpscaledCoordT{UNext[0], UPrev[1], UNext[2], 0})
+                   .template convert<float>() *
+               UCA[0] * (1 - UCA[1]) * UCA[2];
+  auto CA011 = Read(UpscaledCoordT{UPrev[0], UNext[1], UNext[2], 0})
+                   .template convert<float>() *
+               (1 - UCA[0]) * UCA[1] * UCA[2];
+  auto CA111 = Read(UpscaledCoordT{UNext[0], UNext[1], UNext[2], 0})
+                   .template convert<float>() *
+               UCA[0] * UCA[1] * UCA[2];
+  return CA000 + CA100 + CA010 + CA110 + CA001 + CA101 + CA011 + CA111;
+}

sycl/test-e2e/Basic/sycl_2020_images/common.hpp

@@ -14,74 +14,6 @@ constexpr size_t IMAGE_DEPTH = 2;
 constexpr size_t IMAGE_PITCH_WIDTH = 7;
 constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH;
 
-// Implemented as specified by the OpenCL 1.2 specification for
-// CLK_FILTER_LINEAR.
-template <image_format Format, addressing_mode AddrMode, int Dims>
-float4 CalcLinearRead(typename FormatTraits<Format>::rep_elem_type *RefData,
-                      CoordT<ImageType::Sampled, Dims> Coord,
-                      range<2> ImagePitch, range<Dims> ImageRange,
-                      bool Normalized) {
-  using UpscaledCoordT = CoordT<ImageType::Sampled, 3>;
-
-  auto Read = [&](UpscaledCoordT UpCoord) {
-    auto DownCoord = DownscaleCoord<Dims>(UpCoord);
-    return SimulateRead<Format, ImageType::Sampled>(
-        RefData, DownCoord, ImagePitch, ImageRange, false);
-  };
-
-  CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
-  if constexpr (AddrMode == addressing_mode::repeat) {
-    assert(Normalized);
-    AdjCoord -= floor(AdjCoord);
-    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
-  } else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
-    assert(Normalized);
-    AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
-    AdjCoord = sycl::fabs(Coord - AdjCoord);
-    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
-  } else {
-    if (Normalized)
-      AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
-  }
-
-  auto Prev = sycl::floor(AdjCoord - 0.5f);
-  auto Next = Prev + 1;
-  auto CA = (AdjCoord - 0.5f) - Prev;
-
-  Prev = ApplyAddressingMode<AddrMode>(Prev, ImageRange);
-  Next = ApplyAddressingMode<AddrMode>(Next, ImageRange);
-
-  auto UPrev = UpscaleCoord<Dims>(Prev);
-  auto UNext = UpscaleCoord<Dims>(Next);
-  auto UCA = UpscaleCoord<Dims>(CA, 1);
-
-  auto CA000 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UPrev[2], 0})
-                   .template convert<float>() *
-               (1 - UCA[0]) * (1 - UCA[1]) * (1 - UCA[2]);
-  auto CA100 = Read(UpscaledCoordT{UNext[0], UPrev[1], UPrev[2], 0})
-                   .template convert<float>() *
-               UCA[0] * (1 - UCA[1]) * (1 - UCA[2]);
-  auto CA010 = Read(UpscaledCoordT{UPrev[0], UNext[1], UPrev[2], 0})
-                   .template convert<float>() *
-               (1 - UCA[0]) * UCA[1] * (1 - UCA[2]);
-  auto CA110 = Read(UpscaledCoordT{UNext[0], UNext[1], UPrev[2], 0})
-                   .template convert<float>() *
-               UCA[0] * UCA[1] * (1 - UCA[2]);
-  auto CA001 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UNext[2], 0})
-                   .template convert<float>() *
-               (1 - UCA[0]) * (1 - UCA[1]) * UCA[2];
-  auto CA101 = Read(UpscaledCoordT{UNext[0], UPrev[1], UNext[2], 0})
-                   .template convert<float>() *
-               UCA[0] * (1 - UCA[1]) * UCA[2];
-  auto CA011 = Read(UpscaledCoordT{UPrev[0], UNext[1], UNext[2], 0})
-                   .template convert<float>() *
-               (1 - UCA[0]) * UCA[1] * UCA[2];
-  auto CA111 = Read(UpscaledCoordT{UNext[0], UNext[1], UNext[2], 0})
-                   .template convert<float>() *
-               UCA[0] * UCA[1] * UCA[2];
-  return CA000 + CA100 + CA010 + CA110 + CA001 + CA101 + CA011 + CA111;
-}
-
 template <image_format Format, addressing_mode AddrMode,
           coordinate_normalization_mode CoordNormMode, int Dims>
 bool checkSampledImageHostReadLinear(

sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_linear.cpp

@@ -14,35 +14,6 @@ constexpr size_t IMAGE_DEPTH = 2;
 constexpr size_t IMAGE_PITCH_WIDTH = 7;
 constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH;
 
-// Implemented as specified by the OpenCL 1.2 specification for
-// CLK_FILTER_NEAREST.
-template <image_format Format, addressing_mode AddrMode, int Dims>
-typename FormatTraits<Format>::pixel_type
-ReadNearest(typename FormatTraits<Format>::rep_elem_type *RefData,
-            CoordT<ImageType::Sampled, Dims> Coord, range<2> ImagePitch,
-            range<Dims> ImageRange, bool Normalized) {
-  CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
-  if constexpr (AddrMode == addressing_mode::repeat) {
-    assert(Normalized);
-    AdjCoord -= sycl::floor(AdjCoord);
-    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
-    AdjCoord = sycl::floor(AdjCoord);
-  } else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
-    assert(Normalized);
-    AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
-    AdjCoord = sycl::fabs(Coord - AdjCoord);
-    AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
-    AdjCoord = sycl::floor(AdjCoord);
-  } else {
-    if (Normalized)
-      AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
-    AdjCoord = sycl::floor(AdjCoord);
-  }
-  AdjCoord = ApplyAddressingMode<AddrMode>(AdjCoord, ImageRange);
-  return SimulateRead<Format, ImageType::Sampled>(RefData, AdjCoord, ImagePitch,
-                                                  ImageRange, false);
-}
-
 template <image_format Format, addressing_mode AddrMode,
           coordinate_normalization_mode CoordNormMode, int Dims>
 bool checkSampledImageHostReadNearest(