From b92e55dbccc907c23d8aa27855b9e6e3cd961c78 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 1 Jun 2023 05:24:17 -0700 Subject: [PATCH 1/4] [SYCL] Add host_task image accessor support 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 --- sycl/include/sycl/accessor.hpp | 236 +++++++++++++----- sycl/include/sycl/detail/handler_proxy.hpp | 9 +- sycl/include/sycl/ext/oneapi/owner_less.hpp | 12 + sycl/include/sycl/handler.hpp | 13 +- sycl/source/detail/handler_proxy.cpp | 10 + sycl/source/handler.cpp | 26 +- .../Basic/sycl_2020_images/common.hpp | 118 ++++++++- .../host_sampled_image_read_linear.cpp | 68 ----- .../host_sampled_image_read_nearest.cpp | 29 --- .../host_task_sampled_image_read_linear.cpp | 170 +++++++++++++ .../host_task_sampled_image_read_nearest.cpp | 160 ++++++++++++ .../host_task_unsampled_image_read.cpp | 116 +++++++++ .../host_task_unsampled_image_write.cpp | 120 +++++++++ .../host_unsampled_image_write.cpp | 20 +- .../test-e2e/WeakObject/weak_object_utils.hpp | 57 +++++ sycl/test/abi/sycl_symbols_linux.dump | 8 +- 16 files changed, 990 insertions(+), 182 deletions(-) create mode 100644 sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_linear.cpp create mode 100644 sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_nearest.cpp create mode 100644 sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_read.cpp create mode 100644 sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_write.cpp diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 935536378e206..163c5f7dc07f9 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -616,6 +616,41 @@ class __SYCL_EXPORT UnsampledImageAccessorBaseHost { friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj); UnsampledImageAccessorImplPtr impl; + + // The function references helper methods required by GDB pretty-printers + void GDBMethodsAnchor() { +#ifndef NDEBUG + const auto *this_const = this; + (void)getSize(); + (void)this_const->getSize(); + (void)getPtr(); + (void)this_const->getPtr(); +#endif + } + +#ifndef __SYCL_DEVICE_ONLY__ + // Reads a pixel of the underlying image at the specified coordinate. It is + // the responsibility of the caller to ensure that the coordinate type is + // valid. + template + DataT read(const CoordT &Coords) const noexcept { + image_sampler Smpl{addressing_mode::none, + coordinate_normalization_mode::unnormalized, + filtering_mode::nearest}; + return imageReadSamplerHostImpl( + Coords, Smpl, getSize(), getPitch(), getChannelType(), + getChannelOrder(), getPtr(), getElementSize()); + } + + // Writes to a pixel of the underlying image at the specified coordinate. It + // is the responsibility of the caller to ensure that the coordinate type is + // valid. + template + void write(const CoordT &Coords, const DataT &Color) const { + imageWriteHostImpl(Coords, Color, getPitch(), getElementSize(), + getChannelType(), getChannelOrder(), getPtr()); + } +#endif }; class __SYCL_EXPORT SampledImageAccessorBaseHost { @@ -651,6 +686,29 @@ class __SYCL_EXPORT SampledImageAccessorBaseHost { friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj); SampledImageAccessorImplPtr impl; + + // The function references helper methods required by GDB pretty-printers + void GDBMethodsAnchor() { +#ifndef NDEBUG + const auto *this_const = this; + (void)getSize(); + (void)this_const->getSize(); + (void)getPtr(); + (void)this_const->getPtr(); +#endif + } + +#ifndef __SYCL_DEVICE_ONLY__ + // Reads a pixel of the underlying image at the specified coordinate. It is + // the responsibility of the caller to ensure that the coordinate type is + // valid. + template + DataT read(const CoordT &Coords) const { + return imageReadSamplerHostImpl( + Coords, getSampler(), getSize(), getPitch(), getChannelType(), + getChannelOrder(), getPtr(), getElementSize()); + } +#endif }; template struct IsValidCoordDataT; @@ -3504,12 +3562,26 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, template -class __SYCL_EBO unsampled_image_accessor { +class __SYCL_EBO unsampled_image_accessor : +#ifndef __SYCL_DEVICE_ONLY__ + private detail::UnsampledImageAccessorBaseHost, +#endif // __SYCL_DEVICE_ONLY__ + public detail::OwnerLessBase< + unsampled_image_accessor> { static_assert(std::is_same_v || std::is_same_v || std::is_same_v || std::is_same_v, "The data type of an image accessor must be only int4, " "uint4, float4 or half4 from SYCL namespace"); + static_assert(AccessMode == access_mode::read || + AccessMode == access_mode::write, + "Access target must be either read or write."); + +#ifdef __SYCL_DEVICE_ONLY__ + char MPadding[sizeof(detail::UnsampledImageAccessorBaseHost)]; +#else + using host_base_class = detail::UnsampledImageAccessorBaseHost; +#endif // __SYCL_DEVICE_ONLY__ public: using value_type = typename std::conditional unsampled_image_accessor(unsampled_image &ImageRef, handler &CommandGroupHandlerRef, - const property_list &PropList = {}) { + const property_list &PropList = {}) +#ifdef __SYCL_DEVICE_ONLY__ + {} +#else + : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + AccessMode, detail::getSyclObjImpl(ImageRef).get(), + Dimensions, ImageRef.getElementSize(), + {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0}, + ImageRef.getChannelType(), ImageRef.getChannelOrder(), + PropList) { device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef); - if (AccessTarget == image_target::device && !Device.has(aspect::image)) + // Avoid aspect::image warning. + aspect ImageAspect = aspect::image; + if (AccessTarget == image_target::device && !Device.has(ImageAspect)) throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), "Device associated with command group handler does not have " "aspect::image."); - std::ignore = ImageRef; - std::ignore = PropList; - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "unsampled_image_accessor is not yet available."); + detail::associateWithHandler(CommandGroupHandlerRef, this, AccessTarget); + GDBMethodsAnchor(); } +#endif // __SYCL_DEVICE_ONLY__ /* -- common interface members -- */ @@ -3552,10 +3633,7 @@ class __SYCL_EBO unsampled_image_accessor { bool operator==(const unsampled_image_accessor &Rhs) const; #else bool operator==(const unsampled_image_accessor &Rhs) const { - std::ignore = Rhs; - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "operator== is not yet implemented."); + return Rhs.impl == impl; } #endif // __SYCL_DEVICE_ONLY__ @@ -3567,11 +3645,10 @@ class __SYCL_EBO unsampled_image_accessor { size_t size() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ + // Currently not reachable on device. return 0; #else - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "size() is not yet implemented."); + return host_base_class::getSize().size(); #endif // __SYCL_DEVICE_ONLY__ } @@ -3584,13 +3661,12 @@ class __SYCL_EBO unsampled_image_accessor { detail::IsValidUnsampledCoord2020DataT< Dimensions, CoordT>::value>> DataT read(const CoordT &Coords) const noexcept { - std::ignore = Coords; #ifdef __SYCL_DEVICE_ONLY__ + // Currently not reachable on device. + std::ignore = Coords; return {0, 0, 0, 0}; #else - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "read() is not yet implemented."); + return host_base_class::read(Coords); #endif // __SYCL_DEVICE_ONLY__ } @@ -3603,14 +3679,29 @@ class __SYCL_EBO unsampled_image_accessor { detail::IsValidUnsampledCoord2020DataT< Dimensions, CoordT>::value>> void write(const CoordT &Coords, const DataT &Color) const { +#ifdef __SYCL_DEVICE_ONLY__ + // Currently not reachable on device. std::ignore = Coords; std::ignore = Color; +#else + host_base_class::write(Coords, Color); +#endif // __SYCL_DEVICE_ONLY__ + } + +private: + unsampled_image_accessor(const detail::UnsampledImageAccessorImplPtr &Impl) #ifndef __SYCL_DEVICE_ONLY__ - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "write() is not yet implemented."); + : host_base_class{Impl} #endif // __SYCL_DEVICE_ONLY__ + { + std::ignore = Impl; } + + template + friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject); + + template + friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj); }; template ( - Coords, Smpl, base_class::getSize(), base_class::getPitch(), - base_class::getChannelType(), base_class::getChannelOrder(), - base_class::getPtr(), base_class::getElementSize()); + return base_class::read(Coords); } #endif @@ -3720,10 +3805,7 @@ class __SYCL_EBO host_unsampled_image_accessor { // Host implementation is only available in host code. Device is not allowed // to use host_unsampled_image_accessor. - detail::imageWriteHostImpl( - Coords, Color, base_class::getPitch(), base_class::getElementSize(), - base_class::getChannelType(), base_class::getChannelOrder(), - base_class::getPtr()); + base_class::write(Coords, Color); } #endif @@ -3741,13 +3823,24 @@ class __SYCL_EBO host_unsampled_image_accessor template -class __SYCL_EBO sampled_image_accessor { +class __SYCL_EBO sampled_image_accessor : +#ifndef __SYCL_DEVICE_ONLY__ + private detail::SampledImageAccessorBaseHost, +#endif // __SYCL_DEVICE_ONLY__ + public detail::OwnerLessBase< + sampled_image_accessor> { static_assert(std::is_same_v || std::is_same_v || std::is_same_v || std::is_same_v, "The data type of an image accessor must be only int4, " "uint4, float4 or half4 from SYCL namespace"); +#ifdef __SYCL_DEVICE_ONLY__ + char MPadding[sizeof(detail::SampledImageAccessorBaseHost)]; +#else + using host_base_class = detail::SampledImageAccessorBaseHost; +#endif // __SYCL_DEVICE_ONLY__ + public: using value_type = const DataT; using reference = const DataT &; @@ -3756,20 +3849,29 @@ class __SYCL_EBO sampled_image_accessor { template sampled_image_accessor(sampled_image &ImageRef, handler &CommandGroupHandlerRef, - const property_list &PropList = {}) { + const property_list &PropList = {}) +#ifdef __SYCL_DEVICE_ONLY__ + {} +#else + : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + detail::getSyclObjImpl(ImageRef).get(), Dimensions, + ImageRef.getElementSize(), + {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0}, + ImageRef.getChannelType(), ImageRef.getChannelOrder(), + ImageRef.getSampler(), PropList) { device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef); - if (AccessTarget == image_target::device && !Device.has(aspect::image)) + // Avoid aspect::image warning. + aspect ImageAspect = aspect::image; + if (AccessTarget == image_target::device && !Device.has(ImageAspect)) throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), "Device associated with command group handler does not have " "aspect::image."); - std::ignore = ImageRef; - std::ignore = PropList; - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "sampled_image_accessor is not yet available."); + detail::associateWithHandler(CommandGroupHandlerRef, this, AccessTarget); + GDBMethodsAnchor(); } +#endif // __SYCL_DEVICE_ONLY__ /* -- common interface members -- */ @@ -3788,10 +3890,7 @@ class __SYCL_EBO sampled_image_accessor { bool operator==(const sampled_image_accessor &Rhs) const; #else bool operator==(const sampled_image_accessor &Rhs) const { - std::ignore = Rhs; - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "operator== is not yet implemented."); + return Rhs.impl == impl; } #endif // __SYCL_DEVICE_ONLY__ @@ -3803,11 +3902,10 @@ class __SYCL_EBO sampled_image_accessor { size_t size() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ + // Currently not reachable on device. return 0; #else - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "size() is not yet implemented."); + return host_base_class::getSize().size(); #endif // __SYCL_DEVICE_ONLY__ } @@ -3818,15 +3916,29 @@ class __SYCL_EBO sampled_image_accessor { typename = std::enable_if_t::value>> DataT read(const CoordT &Coords) const noexcept { - std::ignore = Coords; #ifdef __SYCL_DEVICE_ONLY__ + // Currently not reachable on device. + std::ignore = Coords; return {0, 0, 0, 0}; #else - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "read() is not yet implemented."); + return host_base_class::read(Coords); +#endif // __SYCL_DEVICE_ONLY__ + } + +private: + sampled_image_accessor(const detail::SampledImageAccessorImplPtr &Impl) +#ifndef __SYCL_DEVICE_ONLY__ + : host_base_class{Impl} #endif // __SYCL_DEVICE_ONLY__ + { + std::ignore = Impl; } + + template + friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject); + + template + friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj); }; template @@ -3897,11 +4009,7 @@ class __SYCL_EBO host_sampled_image_accessor { // Host implementation is only available in host code. Device is not allowed // to use host_sampled_image_accessor. - return detail::imageReadSamplerHostImpl( - Coords, base_class::getSampler(), base_class::getSize(), - base_class::getPitch(), base_class::getChannelType(), - base_class::getChannelOrder(), base_class::getPtr(), - base_class::getElementSize()); + return base_class::read(Coords); } #endif @@ -3984,9 +4092,14 @@ struct hash; size_t operator()(const AccType &A) const { - // TODO: Implement. +#ifdef __SYCL_DEVICE_ONLY__ + // Hash is not supported on DEVICE. Just return 0 here. (void)A; return 0; +#else + auto AccImplPtr = sycl::detail::getSyclObjImpl(A); + return hash()(AccImplPtr); +#endif } }; @@ -4007,9 +4120,14 @@ struct hash> { using AccType = sycl::sampled_image_accessor; size_t operator()(const AccType &A) const { - // TODO: Implement. +#ifdef __SYCL_DEVICE_ONLY__ + // Hash is not supported on DEVICE. Just return 0 here. (void)A; return 0; +#else + auto AccImplPtr = sycl::detail::getSyclObjImpl(A); + return hash()(AccImplPtr); +#endif } }; diff --git a/sycl/include/sycl/detail/handler_proxy.hpp b/sycl/include/sycl/detail/handler_proxy.hpp index c8bb5adefbeaa..a04ca37af7c2a 100644 --- a/sycl/include/sycl/detail/handler_proxy.hpp +++ b/sycl/include/sycl/detail/handler_proxy.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) diff --git a/sycl/include/sycl/ext/oneapi/owner_less.hpp b/sycl/include/sycl/ext/oneapi/owner_less.hpp index 9e09c73a33430..65702d81adabb 100644 --- a/sycl/include/sycl/ext/oneapi/owner_less.hpp +++ b/sycl/include/sycl/ext/oneapi/owner_less.hpp @@ -102,11 +102,23 @@ template struct owner_less> : public detail::owner_less_base> {}; +template +struct owner_less< + unsampled_image_accessor> + : public detail::owner_less_base> {}; + template struct owner_less> : public detail::owner_less_base< host_unsampled_image_accessor> {}; +template +struct owner_less> + : public detail::owner_less_base< + sampled_image_accessor> {}; + template struct owner_less> : public detail::owner_less_base< diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 377fee34ecd89..9d72db7515c22 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.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. @@ -2937,6 +2944,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; diff --git a/sycl/source/detail/handler_proxy.cpp b/sycl/source/detail/handler_proxy.cpp index e3cc8b7c57e70..34d6214922fc5 100644 --- a/sycl/source/detail/handler_proxy.cpp +++ b/sycl/source/detail/handler_proxy.cpp @@ -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 diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 091d125f303c7..5358524548dc2 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -378,9 +378,8 @@ void handler::addReduction(const std::shared_ptr &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); @@ -389,8 +388,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(AccTarget), - /*index*/ 0); + Req, AccTarget, /*index*/ 0); +} + +void handler::associateWithHandler(detail::AccessorBaseHost *AccBase, + access::target AccTarget) { + associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase), + static_cast(AccTarget)); +} + +void handler::associateWithHandler( + detail::UnsampledImageAccessorBaseHost *AccBase, image_target AccTarget) { + associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase), + static_cast(AccTarget)); +} + +void handler::associateWithHandler( + detail::SampledImageAccessorBaseHost *AccBase, image_target AccTarget) { + associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase), + static_cast(AccTarget)); } static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, diff --git a/sycl/test-e2e/Basic/sycl_2020_images/common.hpp b/sycl/test-e2e/Basic/sycl_2020_images/common.hpp index 3d5a65cf92758..0624aa4d82edf 100644 --- a/sycl/test-e2e/Basic/sycl_2020_images/common.hpp +++ b/sycl/test-e2e/Basic/sycl_2020_images/common.hpp @@ -200,7 +200,7 @@ CoordT DelinearizeToCoord(size_t Idx, range ImageRange, } else if constexpr (Dims == 2) { Out = CoordT{Idx % ImageRange[0], Idx / ImageRange[0]}; } else { - Out = CoordT{Idx % ImageRange[0] % ImageRange[1], + Out = CoordT{Idx % ImageRange[0], Idx / ImageRange[0] % ImageRange[1], Idx / ImageRange[0] / ImageRange[1], 0}; } @@ -328,3 +328,119 @@ ApplyAddressingMode(CoordT Coord, } } } + +template static constexpr size_t getMaxInt() { + using rep_elem_type = typename FormatTraits::rep_elem_type; + return static_cast(std::numeric_limits::max()); +} + +template +typename FormatTraits::pixel_type PickNewColor(size_t I, + size_t AccSize) { + using PixelType = typename FormatTraits::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()}); + if constexpr (FormatTraits::Normalized) + NewColor /= AccSize * 4; + return NewColor; +} + +// Implemented as specified by the OpenCL 1.2 specification for +// CLK_FILTER_NEAREST. +template +typename FormatTraits::pixel_type +ReadNearest(typename FormatTraits::rep_elem_type *RefData, + CoordT Coord, range<2> ImagePitch, + range ImageRange, bool Normalized) { + CoordT AdjCoord = Coord; + if constexpr (AddrMode == addressing_mode::repeat) { + assert(Normalized); + AdjCoord -= sycl::floor(AdjCoord); + AdjCoord *= RangeToCoord(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(ImageRange); + AdjCoord = sycl::floor(AdjCoord); + } else { + if (Normalized) + AdjCoord *= RangeToCoord(ImageRange); + AdjCoord = sycl::floor(AdjCoord); + } + AdjCoord = ApplyAddressingMode(AdjCoord, ImageRange); + return SimulateRead(RefData, AdjCoord, ImagePitch, + ImageRange, false); +} + +// Implemented as specified by the OpenCL 1.2 specification for +// CLK_FILTER_LINEAR. +template +float4 CalcLinearRead(typename FormatTraits::rep_elem_type *RefData, + CoordT Coord, + range<2> ImagePitch, range ImageRange, + bool Normalized) { + using UpscaledCoordT = CoordT; + + auto Read = [&](UpscaledCoordT UpCoord) { + auto DownCoord = DownscaleCoord(UpCoord); + return SimulateRead( + RefData, DownCoord, ImagePitch, ImageRange, false); + }; + + CoordT AdjCoord = Coord; + if constexpr (AddrMode == addressing_mode::repeat) { + assert(Normalized); + AdjCoord -= floor(AdjCoord); + AdjCoord *= RangeToCoord(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(ImageRange); + } else { + if (Normalized) + AdjCoord *= RangeToCoord(ImageRange); + } + + auto Prev = sycl::floor(AdjCoord - 0.5f); + auto Next = Prev + 1; + auto CA = (AdjCoord - 0.5f) - Prev; + + Prev = ApplyAddressingMode(Prev, ImageRange); + Next = ApplyAddressingMode(Next, ImageRange); + + auto UPrev = UpscaleCoord(Prev); + auto UNext = UpscaleCoord(Next); + auto UCA = UpscaleCoord(CA, 1); + + auto CA000 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UPrev[2], 0}) + .template convert() * + (1 - UCA[0]) * (1 - UCA[1]) * (1 - UCA[2]); + auto CA100 = Read(UpscaledCoordT{UNext[0], UPrev[1], UPrev[2], 0}) + .template convert() * + UCA[0] * (1 - UCA[1]) * (1 - UCA[2]); + auto CA010 = Read(UpscaledCoordT{UPrev[0], UNext[1], UPrev[2], 0}) + .template convert() * + (1 - UCA[0]) * UCA[1] * (1 - UCA[2]); + auto CA110 = Read(UpscaledCoordT{UNext[0], UNext[1], UPrev[2], 0}) + .template convert() * + UCA[0] * UCA[1] * (1 - UCA[2]); + auto CA001 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UNext[2], 0}) + .template convert() * + (1 - UCA[0]) * (1 - UCA[1]) * UCA[2]; + auto CA101 = Read(UpscaledCoordT{UNext[0], UPrev[1], UNext[2], 0}) + .template convert() * + UCA[0] * (1 - UCA[1]) * UCA[2]; + auto CA011 = Read(UpscaledCoordT{UPrev[0], UNext[1], UNext[2], 0}) + .template convert() * + (1 - UCA[0]) * UCA[1] * UCA[2]; + auto CA111 = Read(UpscaledCoordT{UNext[0], UNext[1], UNext[2], 0}) + .template convert() * + UCA[0] * UCA[1] * UCA[2]; + return CA000 + CA100 + CA010 + CA110 + CA001 + CA101 + CA011 + CA111; +} diff --git a/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_linear.cpp b/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_linear.cpp index 8199dc83862fd..a7be9e774f842 100644 --- a/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_linear.cpp +++ b/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_linear.cpp @@ -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 -float4 CalcLinearRead(typename FormatTraits::rep_elem_type *RefData, - CoordT Coord, - range<2> ImagePitch, range ImageRange, - bool Normalized) { - using UpscaledCoordT = CoordT; - - auto Read = [&](UpscaledCoordT UpCoord) { - auto DownCoord = DownscaleCoord(UpCoord); - return SimulateRead( - RefData, DownCoord, ImagePitch, ImageRange, false); - }; - - CoordT AdjCoord = Coord; - if constexpr (AddrMode == addressing_mode::repeat) { - assert(Normalized); - AdjCoord -= floor(AdjCoord); - AdjCoord *= RangeToCoord(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(ImageRange); - } else { - if (Normalized) - AdjCoord *= RangeToCoord(ImageRange); - } - - auto Prev = sycl::floor(AdjCoord - 0.5f); - auto Next = Prev + 1; - auto CA = (AdjCoord - 0.5f) - Prev; - - Prev = ApplyAddressingMode(Prev, ImageRange); - Next = ApplyAddressingMode(Next, ImageRange); - - auto UPrev = UpscaleCoord(Prev); - auto UNext = UpscaleCoord(Next); - auto UCA = UpscaleCoord(CA, 1); - - auto CA000 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UPrev[2], 0}) - .template convert() * - (1 - UCA[0]) * (1 - UCA[1]) * (1 - UCA[2]); - auto CA100 = Read(UpscaledCoordT{UNext[0], UPrev[1], UPrev[2], 0}) - .template convert() * - UCA[0] * (1 - UCA[1]) * (1 - UCA[2]); - auto CA010 = Read(UpscaledCoordT{UPrev[0], UNext[1], UPrev[2], 0}) - .template convert() * - (1 - UCA[0]) * UCA[1] * (1 - UCA[2]); - auto CA110 = Read(UpscaledCoordT{UNext[0], UNext[1], UPrev[2], 0}) - .template convert() * - UCA[0] * UCA[1] * (1 - UCA[2]); - auto CA001 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UNext[2], 0}) - .template convert() * - (1 - UCA[0]) * (1 - UCA[1]) * UCA[2]; - auto CA101 = Read(UpscaledCoordT{UNext[0], UPrev[1], UNext[2], 0}) - .template convert() * - UCA[0] * (1 - UCA[1]) * UCA[2]; - auto CA011 = Read(UpscaledCoordT{UPrev[0], UNext[1], UNext[2], 0}) - .template convert() * - (1 - UCA[0]) * UCA[1] * UCA[2]; - auto CA111 = Read(UpscaledCoordT{UNext[0], UNext[1], UNext[2], 0}) - .template convert() * - UCA[0] * UCA[1] * UCA[2]; - return CA000 + CA100 + CA010 + CA110 + CA001 + CA101 + CA011 + CA111; -} - template bool checkSampledImageHostReadLinear( diff --git a/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_nearest.cpp b/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_nearest.cpp index 849683979e1b9..afcfbb67d4544 100644 --- a/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_nearest.cpp +++ b/sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_nearest.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 -typename FormatTraits::pixel_type -ReadNearest(typename FormatTraits::rep_elem_type *RefData, - CoordT Coord, range<2> ImagePitch, - range ImageRange, bool Normalized) { - CoordT AdjCoord = Coord; - if constexpr (AddrMode == addressing_mode::repeat) { - assert(Normalized); - AdjCoord -= sycl::floor(AdjCoord); - AdjCoord *= RangeToCoord(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(ImageRange); - AdjCoord = sycl::floor(AdjCoord); - } else { - if (Normalized) - AdjCoord *= RangeToCoord(ImageRange); - AdjCoord = sycl::floor(AdjCoord); - } - AdjCoord = ApplyAddressingMode(AdjCoord, ImageRange); - return SimulateRead(RefData, AdjCoord, ImagePitch, - ImageRange, false); -} - template bool checkSampledImageHostReadNearest( diff --git a/sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_linear.cpp b/sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_linear.cpp new file mode 100644 index 0000000000000..bb87ffeb443a4 --- /dev/null +++ b/sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_linear.cpp @@ -0,0 +1,170 @@ +// RUN: %{build} -o %t.out +// TODO: Consider moving to sycl/test as this is device-independent. +// RUN: %{run-unfiltered-devices} %t.out + +// Tests for indirect read of sampled_image using host accessors and linear +// filtering mode. + +#include "common.hpp" + +constexpr size_t IMAGE_WIDTH = 5; +constexpr size_t IMAGE_HEIGHT = 4; +constexpr size_t IMAGE_DEPTH = 2; + +constexpr size_t IMAGE_PITCH_WIDTH = 7; +constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH; + +template +bool checkSampledImageHostTaskReadLinear( + sampled_image &Image, + typename FormatTraits::rep_elem_type *RefData, queue &Q) { + using PixelType = typename FormatTraits::pixel_type; + constexpr ImageType ImgType = ImageType::Sampled; + constexpr bool Normalized = + CoordNormMode == coordinate_normalization_mode::normalized; + + bool Success = true; + { + buffer SuccessBuf{&Success, 1}; + Q.submit([&](handler &CGH) { + sampled_image_accessor Acc( + Image, CGH); + accessor SuccessAcc{SuccessBuf, CGH, write_only_host_task}; + + assert(Image.size() == Acc.size()); + CGH.host_task([=]() { + range ImageRange = Image.get_range(); + range<2> ImagePitch = getElementWisePitch(Image); + + // Get the offset permutations to add to the direct coordinates to check + // non-direct reads. + auto Offsets = GetOffsetPermutations(); + + for (size_t I = 0; I < Acc.size(); ++I) { + CoordT Coord = + DelinearizeToCoord(I, ImageRange, Normalized); + + for (const auto &Offset : Offsets) { + // Normalize offset if needed. + auto AdjOffset = + Normalized ? Offset / RangeToCoord(ImageRange, 2) + : Offset; + auto OffsetCoord = Coord + AdjOffset; + + float4 ReadVal = Acc.read(OffsetCoord).template convert(); + float4 ExpectedVal = CalcLinearRead( + RefData, OffsetCoord, ImagePitch, ImageRange, Normalized); + + // Compare results as floats to better check for precision + // differences. I.e. if the type is already represented as float we + // can simply use a small precision, while if they are integral the + // precision could have affected the rounding. + float Precision = + std::is_integral_v::rep_elem_type> + ? 0.5001 + : 0.1; + if (!ApproxEq(ReadVal, ExpectedVal, Precision)) { + std::cout << "Unexpected read value (" << ReadVal + << " != " << ExpectedVal << ") at coordinate " + << OffsetCoord << " (" << FormatTraits::Name + << ") (" << AddressingModeToString() << ")" + << std::endl; + SuccessAcc[0] = false; + } + } + } + }); + }); + } + return Success; +} + +template +int check(std::vector::rep_elem_type> &Data, + queue &Q) { + range ImageRange = + CreateImageRange(IMAGE_WIDTH, IMAGE_HEIGHT, IMAGE_DEPTH); + + constexpr image_sampler Sampler{AddrMode, CoordNormMode, + filtering_mode::linear}; + + int Failures = 0; + + // Test image without explicit pitch. + sampled_image Img1{Data.data(), Format, Sampler, ImageRange}; + Failures += + !checkSampledImageHostTaskReadLinear( + Img1, Data.data(), Q); + + // If Dims > 1 test image image with a pitch different than the image size. + if constexpr (Dims > 1) { + constexpr size_t REP_ELEM_VEC_SIZE = + sizeof(typename FormatTraits::rep_elem_type) * 4; + constexpr size_t IMAGE_PITCH_WIDTH_BYTES = + IMAGE_PITCH_WIDTH * REP_ELEM_VEC_SIZE; + constexpr size_t IMAGE_PITCH_HEIGHT_BYTES = + IMAGE_PITCH_HEIGHT * REP_ELEM_VEC_SIZE; + range ImagePitch = CreateImageRange( + IMAGE_PITCH_WIDTH_BYTES, IMAGE_PITCH_HEIGHT_BYTES, 0); + + sampled_image Img2{Data.data(), Format, Sampler, ImageRange, + ImagePitch}; + Failures += + !checkSampledImageHostTaskReadLinear( + Img2, Data.data(), Q); + } + + return Failures; +} + +template +int checkForFormatAndDims( + std::vector::rep_elem_type> &Data, queue &Q) { + int Failures = 0; + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + return Failures; +} + +template int checkForFormat(queue &Q) { + auto Data = GenerateData(IMAGE_PITCH_WIDTH * IMAGE_PITCH_HEIGHT * + IMAGE_DEPTH); + int Failures = 0; + Failures += checkForFormatAndDims(Data, Q); + Failures += checkForFormatAndDims(Data, Q); + Failures += checkForFormatAndDims(Data, Q); + return Failures; +} + +int main() { + queue Q; + int Failures = 0; + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + return Failures; +} diff --git a/sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_nearest.cpp b/sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_nearest.cpp new file mode 100644 index 0000000000000..d69f0db910bfb --- /dev/null +++ b/sycl/test-e2e/Basic/sycl_2020_images/host_task_sampled_image_read_nearest.cpp @@ -0,0 +1,160 @@ +// RUN: %{build} -o %t.out +// TODO: Consider moving to sycl/test as this is device-independent. +// RUN: %{run-unfiltered-devices} %t.out + +// Tests for indirect read of sampled_image using host accessors and linear +// filtering mode. + +#include "common.hpp" + +constexpr size_t IMAGE_WIDTH = 5; +constexpr size_t IMAGE_HEIGHT = 4; +constexpr size_t IMAGE_DEPTH = 2; + +constexpr size_t IMAGE_PITCH_WIDTH = 7; +constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH; + +template +bool checkSampledImageHostTaskReadNearest( + sampled_image &Image, + typename FormatTraits::rep_elem_type *RefData, queue &Q) { + using PixelType = typename FormatTraits::pixel_type; + constexpr ImageType ImgType = ImageType::Sampled; + constexpr bool Normalized = + CoordNormMode == coordinate_normalization_mode::normalized; + + bool Success = true; + { + buffer SuccessBuf{&Success, 1}; + Q.submit([&](handler &CGH) { + sampled_image_accessor Acc( + Image, CGH); + accessor SuccessAcc{SuccessBuf, CGH, write_only_host_task}; + + assert(Image.size() == Acc.size()); + + CGH.host_task([=]() { + range ImageRange = Image.get_range(); + range<2> ImagePitch = getElementWisePitch(Image); + + auto Offsets = GetOffsetPermutations(); + + for (size_t I = 0; I < Acc.size(); ++I) { + CoordT Coord = + DelinearizeToCoord(I, ImageRange, Normalized); + + for (const auto &Offset : Offsets) { + // Normalize offset if needed. + auto AdjOffset = + Normalized ? Offset / RangeToCoord(ImageRange, 2) + : Offset; + auto OffsetCoord = Coord + AdjOffset; + + auto ReadVal = Acc.read(OffsetCoord); + auto ExpectedVal = ReadNearest( + RefData, OffsetCoord, ImagePitch, ImageRange, Normalized); + if (!ApproxEq(ReadVal, ExpectedVal)) { + std::cout << "Unexpected read value (" << ReadVal + << " != " << ExpectedVal << ") at coordinate " + << OffsetCoord << " (" << FormatTraits::Name + << ") (" << AddressingModeToString() << ")" + << std::endl; + SuccessAcc[0] = false; + } + } + } + }); + }); + } + return Success; +} + +template +int check(std::vector::rep_elem_type> &Data, + queue &Q) { + range ImageRange = + CreateImageRange(IMAGE_WIDTH, IMAGE_HEIGHT, IMAGE_DEPTH); + + constexpr image_sampler Sampler{AddrMode, CoordNormMode, + filtering_mode::nearest}; + + int Failures = 0; + + // Test image without explicit pitch. + sampled_image Img1{Data.data(), Format, Sampler, ImageRange}; + Failures += + !checkSampledImageHostTaskReadNearest( + Img1, Data.data(), Q); + + // If Dims > 1 test image image with a pitch different than the image size. + if constexpr (Dims > 1) { + constexpr size_t REP_ELEM_VEC_SIZE = + sizeof(typename FormatTraits::rep_elem_type) * 4; + constexpr size_t IMAGE_PITCH_WIDTH_BYTES = + IMAGE_PITCH_WIDTH * REP_ELEM_VEC_SIZE; + constexpr size_t IMAGE_PITCH_HEIGHT_BYTES = + IMAGE_PITCH_HEIGHT * REP_ELEM_VEC_SIZE; + range ImagePitch = CreateImageRange( + IMAGE_PITCH_WIDTH_BYTES, IMAGE_PITCH_HEIGHT_BYTES, 0); + + sampled_image Img2{Data.data(), Format, Sampler, ImageRange, + ImagePitch}; + Failures += + !checkSampledImageHostTaskReadNearest( + Img2, Data.data(), Q); + } + + return Failures; +} + +template +int checkForFormatAndDims( + std::vector::rep_elem_type> &Data, queue &Q) { + int Failures = 0; + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + return Failures; +} + +template int checkForFormat(queue &Q) { + auto Data = GenerateData(IMAGE_PITCH_WIDTH * IMAGE_PITCH_HEIGHT * + IMAGE_DEPTH); + int Failures = 0; + Failures += checkForFormatAndDims(Data, Q); + Failures += checkForFormatAndDims(Data, Q); + Failures += checkForFormatAndDims(Data, Q); + return Failures; +} + +int main() { + queue Q; + int Failures = 0; + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + return Failures; +} diff --git a/sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_read.cpp b/sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_read.cpp new file mode 100644 index 0000000000000..99774f29cc832 --- /dev/null +++ b/sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_read.cpp @@ -0,0 +1,116 @@ +// RUN: %{build} -o %t.out +// TODO: Consider moving to sycl/test as this is device-independent. +// RUN: %{run-unfiltered-devices} %t.out + +// Tests for direct read of unsampled_image using host accessors. + +#include "common.hpp" + +constexpr size_t IMAGE_WIDTH = 5; +constexpr size_t IMAGE_HEIGHT = 4; +constexpr size_t IMAGE_DEPTH = 2; + +constexpr size_t IMAGE_PITCH_WIDTH = 7; +constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH; + +template +bool checkUnsampledImageHostTaskReadDirect( + unsampled_image &Image, + typename FormatTraits::rep_elem_type *RefData, queue &Q) { + using PixelType = typename FormatTraits::pixel_type; + constexpr ImageType ImgType = ImageType::Unsampled; + + bool Success = true; + { + buffer SuccessBuf{&Success, 1}; + Q.submit([&](handler &CGH) { + unsampled_image_accessor + Acc(Image, CGH); + accessor SuccessAcc{SuccessBuf, CGH, write_only_host_task}; + + assert(Image.size() == Acc.size()); + + CGH.host_task([=]() { + range ImageRange = Image.get_range(); + range<2> ImagePitch = getElementWisePitch(Image); + + for (size_t I = 0; I < Acc.size(); ++I) { + auto Coord = DelinearizeToCoord(I, ImageRange); + + // Read the coordinate through the accessor and read the corresponding + // value in the reference memory. + PixelType ReadVal = Acc.read(Coord); + PixelType ExpectedVal = SimulateRead( + RefData, Coord, ImagePitch, ImageRange); + if (!AllTrue(ReadVal == ExpectedVal)) { + std::cout << "Unexpected read value (" << ReadVal + << " != " << ExpectedVal << ") at coordinate " << Coord + << " (" << FormatTraits::Name << ")" << std::endl; + SuccessAcc[0] = false; + } + } + }); + }); + } + return Success; +} + +template +int check(std::vector::rep_elem_type> &Data, + queue &Q) { + range ImageRange = + CreateImageRange(IMAGE_WIDTH, IMAGE_HEIGHT, IMAGE_DEPTH); + + int Failures = 0; + + // Test image without explicit pitch. + unsampled_image Img1{Data.data(), Format, ImageRange}; + Failures += + !checkUnsampledImageHostTaskReadDirect(Img1, Data.data(), Q); + + // If Dims > 1 test image image with a pitch different than the image size. + if constexpr (Dims > 1) { + constexpr size_t REP_ELEM_VEC_SIZE = + sizeof(typename FormatTraits::rep_elem_type) * 4; + constexpr size_t IMAGE_PITCH_WIDTH_BYTES = + IMAGE_PITCH_WIDTH * REP_ELEM_VEC_SIZE; + constexpr size_t IMAGE_PITCH_HEIGHT_BYTES = + IMAGE_PITCH_HEIGHT * REP_ELEM_VEC_SIZE; + range ImagePitch = CreateImageRange( + IMAGE_PITCH_WIDTH_BYTES, IMAGE_PITCH_HEIGHT_BYTES, 0); + + unsampled_image Img2{Data.data(), Format, ImageRange, ImagePitch}; + Failures += + !checkUnsampledImageHostTaskReadDirect(Img2, Data.data(), Q); + } + + return Failures; +} + +template int checkForFormat(queue &Q) { + auto Data = GenerateData(IMAGE_PITCH_WIDTH * IMAGE_PITCH_HEIGHT * + IMAGE_DEPTH); + int Failures = 0; + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + return Failures; +} + +int main() { + queue Q; + int Failures = 0; + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + return Failures; +} diff --git a/sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_write.cpp b/sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_write.cpp new file mode 100644 index 0000000000000..2fd1f31026e15 --- /dev/null +++ b/sycl/test-e2e/Basic/sycl_2020_images/host_task_unsampled_image_write.cpp @@ -0,0 +1,120 @@ +// RUN: %{build} -o %t.out +// TODO: Consider moving to sycl/test as this is device-independent. +// RUN: %{run-unfiltered-devices} %t.out + +// Tests for direct read of unsampled_image using host accessors. + +#include "common.hpp" + +constexpr size_t IMAGE_WIDTH = 5; +constexpr size_t IMAGE_HEIGHT = 4; +constexpr size_t IMAGE_DEPTH = 2; + +constexpr size_t IMAGE_PITCH_WIDTH = 7; +constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH; + +template +bool checkUnsampledImageHostWriteDirect(unsampled_image &Image, + queue &Q) { + using PixelType = typename FormatTraits::pixel_type; + constexpr ImageType ImgType = ImageType::Unsampled; + + range ImageRange = Image.get_range(); + size_t ImageSize = Image.size(); + + Q.submit([&](handler &CGH) { + unsampled_image_accessor + Acc(Image, CGH); + + CGH.host_task([=]() { + assert(ImageSize == Acc.size()); + + for (size_t I = 0; I < ImageSize; ++I) { + auto Coord = DelinearizeToCoord(I, ImageRange); + PixelType NewColor = PickNewColor(I, ImageSize); + Acc.write(Coord, NewColor); + } + }); + }); + + // Check result on host. + host_unsampled_image_accessor Acc(Image); + + assert(ImageSize == Acc.size()); + + bool success = true; + for (size_t I = 0; I < ImageSize; ++I) { + auto Coord = DelinearizeToCoord(I, ImageRange); + PixelType ExpectedColor = PickNewColor(I, ImageSize); + + PixelType ReadVal = Acc.read(Coord); + if (!ApproxEq(ReadVal, ExpectedColor)) { + std::cout << "Unexpected written value (" << ReadVal + << " != " << ExpectedColor << ") at coordinate " << Coord + << " (" << FormatTraits::Name << ")" << std::endl; + success = false; + } + } + return success; +} + +template +int check(std::vector::rep_elem_type> &Data, + queue &Q) { + range ImageRange = + CreateImageRange(IMAGE_WIDTH, IMAGE_HEIGHT, IMAGE_DEPTH); + + int Failures = 0; + + // Test image without explicit pitch. + std::fill(Data.begin(), Data.end(), 0); + unsampled_image Img1{Data.data(), Format, ImageRange}; + Failures += !checkUnsampledImageHostWriteDirect(Img1, Q); + + // If Dims > 1 test image image with a pitch different than the image size. + if constexpr (Dims > 1) { + constexpr size_t REP_ELEM_VEC_SIZE = + sizeof(typename FormatTraits::rep_elem_type) * 4; + constexpr size_t IMAGE_PITCH_WIDTH_BYTES = + IMAGE_PITCH_WIDTH * REP_ELEM_VEC_SIZE; + constexpr size_t IMAGE_PITCH_HEIGHT_BYTES = + IMAGE_PITCH_HEIGHT * REP_ELEM_VEC_SIZE; + range ImagePitch = CreateImageRange( + IMAGE_PITCH_WIDTH_BYTES, IMAGE_PITCH_HEIGHT_BYTES, 0); + + std::fill(Data.begin(), Data.end(), 0); + unsampled_image Img2{Data.data(), Format, ImageRange, ImagePitch}; + Failures += !checkUnsampledImageHostWriteDirect(Img2, Q); + } + + return Failures; +} + +template int checkForFormat(queue &Q) { + + auto Data = GenerateData(IMAGE_PITCH_WIDTH * IMAGE_PITCH_HEIGHT * + IMAGE_DEPTH); + int Failures = 0; + Failures += check(Data, Q); + Failures += check(Data, Q); + Failures += check(Data, Q); + return Failures; +} + +int main() { + queue Q; + int Failures = 0; + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + Failures += checkForFormat(Q); + return Failures; +} diff --git a/sycl/test-e2e/Basic/sycl_2020_images/host_unsampled_image_write.cpp b/sycl/test-e2e/Basic/sycl_2020_images/host_unsampled_image_write.cpp index 0e03ea5cbd7b3..ce3d514766b60 100644 --- a/sycl/test-e2e/Basic/sycl_2020_images/host_unsampled_image_write.cpp +++ b/sycl/test-e2e/Basic/sycl_2020_images/host_unsampled_image_write.cpp @@ -13,13 +13,8 @@ constexpr size_t IMAGE_DEPTH = 2; constexpr size_t IMAGE_PITCH_WIDTH = 7; constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH; -template static constexpr size_t getMaxInt() { - using rep_elem_type = typename FormatTraits::rep_elem_type; - return static_cast(std::numeric_limits::max()); -} - template -bool checkUnsampledImageHostWriteDirect(unsampled_image &Image) { +bool checkUnsampledImageHostTaskWriteDirect(unsampled_image &Image) { using PixelType = typename FormatTraits::pixel_type; constexpr ImageType ImgType = ImageType::Unsampled; @@ -29,18 +24,11 @@ bool checkUnsampledImageHostWriteDirect(unsampled_image &Image) { assert(Image.size() == AccSize); range ImageRange = Image.get_range(); - range<2> ImagePitch = getElementWisePitch(Image); bool success = true; for (size_t I = 0; I < AccSize; ++I) { auto Coord = DelinearizeToCoord(I, ImageRange); - 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()}); - if constexpr (FormatTraits::Normalized) - NewColor /= AccSize * 4; + PixelType NewColor = PickNewColor(I, AccSize); // Write the new color to the coordinate, then try to read it and check that // it has changed accordingly. @@ -66,7 +54,7 @@ int check(std::vector::rep_elem_type> &Data) { // Test image without explicit pitch. std::fill(Data.begin(), Data.end(), 0); unsampled_image Img1{Data.data(), Format, ImageRange}; - Failures += !checkUnsampledImageHostWriteDirect(Img1); + Failures += !checkUnsampledImageHostTaskWriteDirect(Img1); // If Dims > 1 test image image with a pitch different than the image size. if constexpr (Dims > 1) { @@ -81,7 +69,7 @@ int check(std::vector::rep_elem_type> &Data) { std::fill(Data.begin(), Data.end(), 0); unsampled_image Img2{Data.data(), Format, ImageRange, ImagePitch}; - Failures += !checkUnsampledImageHostWriteDirect(Img2); + Failures += !checkUnsampledImageHostTaskWriteDirect(Img2); } return Failures; diff --git a/sycl/test-e2e/WeakObject/weak_object_utils.hpp b/sycl/test-e2e/WeakObject/weak_object_utils.hpp index fd5e4d8c617e9..030650b423223 100644 --- a/sycl/test-e2e/WeakObject/weak_object_utils.hpp +++ b/sycl/test-e2e/WeakObject/weak_object_utils.hpp @@ -105,6 +105,21 @@ template