From 34f3dfc984e317eeb1fbdbbec84da01c44e8c2ab Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 15 Sep 2025 16:10:04 +0200 Subject: [PATCH 1/3] [SYCL][NFC] Don't use `std::ignore` Its implementation involve templates, which means "extra compilation time overhead". `[[maybe_unused]]` does the job just fine. This patch only changed `sycl/` headers, ignoring `sycl/ext` headers. --- sycl/include/sycl/accessor_image.hpp | 19 ++++++++----------- sycl/include/sycl/detail/spirv.hpp | 8 ++++---- sycl/include/sycl/handler.hpp | 4 +--- sycl/include/sycl/reduction.hpp | 9 ++------- sycl/include/sycl/sub_group.hpp | 11 ++++------- 5 files changed, 19 insertions(+), 32 deletions(-) diff --git a/sycl/include/sycl/accessor_image.hpp b/sycl/include/sycl/accessor_image.hpp index 468bb8d62dacd..a15cf376879b7 100644 --- a/sycl/include/sycl/accessor_image.hpp +++ b/sycl/include/sycl/accessor_image.hpp @@ -904,10 +904,9 @@ class __SYCL_EBO unsampled_image_accessor : typename = std::enable_if_t::value>> - DataT read(const CoordT &Coords) const noexcept { + DataT read([[maybe_unused]] const CoordT &Coords) const noexcept { #ifdef __SYCL_DEVICE_ONLY__ // Currently not reachable on device. - std::ignore = Coords; return {0, 0, 0, 0}; #else return host_base_class::read(Coords); @@ -922,23 +921,22 @@ class __SYCL_EBO unsampled_image_accessor : typename = std::enable_if_t::value>> - void write(const CoordT &Coords, const DataT &Color) const { + void write([[maybe_unused]] const CoordT &Coords, + [[maybe_unused]] 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) + unsampled_image_accessor( + [[maybe_unused]] const detail::UnsampledImageAccessorImplPtr &Impl) #ifndef __SYCL_DEVICE_ONLY__ : host_base_class{Impl} #endif // __SYCL_DEVICE_ONLY__ { - std::ignore = Impl; } template @@ -1213,10 +1211,9 @@ class __SYCL_EBO sampled_image_accessor : template ::value>> - DataT read(const CoordT &Coords) const noexcept { + DataT read([[maybe_unused]] const CoordT &Coords) const noexcept { #ifdef __SYCL_DEVICE_ONLY__ // Currently not reachable on device. - std::ignore = Coords; return {0, 0, 0, 0}; #else return host_base_class::read(Coords); @@ -1224,12 +1221,12 @@ class __SYCL_EBO sampled_image_accessor : } private: - sampled_image_accessor(const detail::SampledImageAccessorImplPtr &Impl) + sampled_image_accessor( + [[maybe_unused]] const detail::SampledImageAccessorImplPtr &Impl) #ifndef __SYCL_DEVICE_ONLY__ : host_base_class{Impl} #endif // __SYCL_DEVICE_ONLY__ { - std::ignore = Impl; } template diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp index 671e4758a268d..a62ad798d35f9 100644 --- a/sycl/include/sycl/detail/spirv.hpp +++ b/sycl/include/sycl/detail/spirv.hpp @@ -921,10 +921,10 @@ template EnableIfGenericShuffle ShuffleUp(GroupT g, T x, uint32_t delta); template -EnableIfNativeShuffle Shuffle(GroupT g, T x, id<1> local_id) { +EnableIfNativeShuffle Shuffle([[maybe_unused]] GroupT g, T x, + id<1> local_id) { uint32_t LocalId = MapShuffleID(g, local_id); #ifndef __NVPTX__ - std::ignore = g; if constexpr (ext::oneapi::experimental::is_user_constructed_group_v< GroupT> && detail::is_vec::value) { @@ -955,9 +955,9 @@ EnableIfNativeShuffle Shuffle(GroupT g, T x, id<1> local_id) { } template -EnableIfNativeShuffle ShuffleXor(GroupT g, T x, id<1> mask) { +EnableIfNativeShuffle ShuffleXor([[maybe_unused]] GroupT g, T x, + id<1> mask) { #ifndef __NVPTX__ - std::ignore = g; if constexpr (ext::oneapi::experimental::is_user_constructed_group_v< GroupT> && detail::is_vec::value) { diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index a70cc50d5f42c..0176ee8a62d30 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -902,7 +902,7 @@ class __SYCL_EXPORT handler { /// /// Stores information about kernel properties into the handler. template - void processLaunchProperties(PropertiesT Props) { + void processLaunchProperties([[maybe_unused]] PropertiesT Props) { if constexpr (PropertiesT::template has_property< sycl::ext::intel::experimental::cache_config_key>()) { auto Config = Props.template get_property< @@ -912,8 +912,6 @@ class __SYCL_EXPORT handler { } else if (Config == sycl::ext::intel::experimental::large_data) { setKernelCacheConfig(StableKernelCacheConfig::LargeData); } - } else { - std::ignore = Props; } constexpr bool UsesRootSync = PropertiesT::template has_property< diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp index d7afe62562ff3..0274651d7f50b 100644 --- a/sycl/include/sycl/reduction.hpp +++ b/sycl/include/sycl/reduction.hpp @@ -1036,7 +1036,6 @@ class reduction_impl_algo { template std::enable_if_t withInitializedMem(handler &CGH, FuncTy Func) { - std::ignore = CGH; assert(!initializeToIdentity() && "Initialize to identity not allowed for identity-less reductions."); Func(accessor{MRedOut, CGH}); @@ -1079,7 +1078,6 @@ class reduction_impl_algo { bool initializeToIdentity() const { return InitializeToIdentity; } auto getUserRedVarAccess(handler &CGH) { - std::ignore = CGH; if constexpr (is_usm) return MRedOut; else @@ -1816,7 +1814,6 @@ template <> struct NDRangeReduction { size_t NWorkGroups = NDRange.get_group_range().size(); bool IsUpdateOfUserVar = !Redu.initializeToIdentity(); - std::ignore = IsUpdateOfUserVar; // The type of the Out "accessor" differs between scenarios when there is // just one WorkGroup and when there are multiple. Use this lambda to write @@ -2800,9 +2797,8 @@ void reduction_parallel_for(handler &CGH, range Range, /// The reduction algorithm may be less efficient if the specified binary /// operation does not have a known identity. template -auto reduction(buffer Var, handler &CGH, +auto reduction(buffer Var, handler &, BinaryOperation Combiner, const property_list &PropList = {}) { - std::ignore = CGH; detail::verifyReductionProps(PropList); bool InitializeToIdentity = PropList.has_property(); @@ -2829,9 +2825,8 @@ auto reduction(T *Var, BinaryOperation Combiner, /// reduction identity value \p Identity, reduction operation \p Combiner, /// and optional reduction properties. template -auto reduction(buffer Var, handler &CGH, const T &Identity, +auto reduction(buffer Var, handler &, const T &Identity, BinaryOperation Combiner, const property_list &PropList = {}) { - std::ignore = CGH; detail::verifyReductionProps(PropList); bool InitializeToIdentity = PropList.has_property(); diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp index 4bad8036f5c63..f053f1ca994a0 100644 --- a/sycl/include/sycl/sub_group.hpp +++ b/sycl/include/sycl/sub_group.hpp @@ -21,7 +21,6 @@ #include // for range #include // for uint32_t -#include // for _Swallow_assign, ignore #include // for enable_if_t, remove_cv_t namespace sycl { @@ -633,23 +632,21 @@ struct sub_group { } // Common member functions for by-value semantics - friend bool operator==(const sub_group &lhs, const sub_group &rhs) { + friend bool operator==([[maybe_unused]] const sub_group &lhs, + [[maybe_unused]] const sub_group &rhs) { #ifdef __SYCL_DEVICE_ONLY__ return lhs.get_group_id() == rhs.get_group_id(); #else - std::ignore = lhs; - std::ignore = rhs; throw sycl::exception(make_error_code(errc::feature_not_supported), "Sub-groups are not supported on host."); #endif } - friend bool operator!=(const sub_group &lhs, const sub_group &rhs) { + friend bool operator!=([[maybe_unused]] const sub_group &lhs, + [[maybe_unused]] const sub_group &rhs) { #ifdef __SYCL_DEVICE_ONLY__ return !(lhs == rhs); #else - std::ignore = lhs; - std::ignore = rhs; throw sycl::exception(make_error_code(errc::feature_not_supported), "Sub-groups are not supported on host."); #endif From 9b44a1d6d6c7b74fd0bbfec9ad6356c3f5f18f68 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 15 Sep 2025 17:42:32 +0200 Subject: [PATCH 2/3] This should presumably fix gcc 7.5 build --- sycl/include/sycl/accessor_image.hpp | 146 +----------------- sycl/include/sycl/accessor_image_base.hpp | 172 ++++++++++++++++++++++ sycl/source/detail/accessor_impl.cpp | 10 +- sycl/source/detail/accessor_impl.hpp | 2 +- sycl/source/handler.cpp | 1 + 5 files changed, 181 insertions(+), 150 deletions(-) create mode 100644 sycl/include/sycl/accessor_image_base.hpp diff --git a/sycl/include/sycl/accessor_image.hpp b/sycl/include/sycl/accessor_image.hpp index a15cf376879b7..235006ec8d2f1 100644 --- a/sycl/include/sycl/accessor_image.hpp +++ b/sycl/include/sycl/accessor_image.hpp @@ -1,6 +1,7 @@ #pragma once #include +#include #include #include #include @@ -61,156 +62,11 @@ void __SYCL_EXPORT sampledImageConstructorNotification( const std::optional &Target, const void *Type, uint32_t ElemSize, const code_location &CodeLoc); -class UnsampledImageAccessorImplHost; -class SampledImageAccessorImplHost; -using UnsampledImageAccessorImplPtr = - std::shared_ptr; -using SampledImageAccessorImplPtr = - std::shared_ptr; - void __SYCL_EXPORT addHostUnsampledImageAccessorAndWait(UnsampledImageAccessorImplHost *Req); void __SYCL_EXPORT addHostSampledImageAccessorAndWait(SampledImageAccessorImplHost *Req); -class __SYCL_EXPORT UnsampledImageAccessorBaseHost { -protected: - UnsampledImageAccessorBaseHost(const UnsampledImageAccessorImplPtr &Impl) - : impl{Impl} {} - -public: - UnsampledImageAccessorBaseHost(sycl::range<3> Size, access_mode AccessMode, - void *SYCLMemObject, int Dims, int ElemSize, - id<3> Pitch, image_channel_type ChannelType, - image_channel_order ChannelOrder, - const property_list &PropertyList = {}); - const sycl::range<3> &getSize() const; - void *getMemoryObject() const; - detail::AccHostDataT &getAccData(); - void *getPtr(); - void *getPtr() const; - int getNumOfDims() const; - int getElementSize() const; - id<3> getPitch() const; - image_channel_type getChannelType() const; - image_channel_order getChannelOrder() const; - const property_list &getPropList() const; - -protected: - template - friend const decltype(Obj::impl) & - detail::getSyclObjImpl(const Obj &SyclObject); - - template - friend T detail::createSyclObjFromImpl( - std::add_rvalue_reference_t ImplObj); - - template - friend T detail::createSyclObjFromImpl( - std::add_lvalue_reference_t 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 { -protected: - SampledImageAccessorBaseHost(const SampledImageAccessorImplPtr &Impl) - : impl{Impl} {} - -public: - SampledImageAccessorBaseHost(sycl::range<3> Size, void *SYCLMemObject, - int Dims, int ElemSize, id<3> Pitch, - image_channel_type ChannelType, - image_channel_order ChannelOrder, - image_sampler Sampler, - const property_list &PropertyList = {}); - const sycl::range<3> &getSize() const; - void *getMemoryObject() const; - detail::AccHostDataT &getAccData(); - void *getPtr(); - void *getPtr() const; - int getNumOfDims() const; - int getElementSize() const; - id<3> getPitch() const; - image_channel_type getChannelType() const; - image_channel_order getChannelOrder() const; - image_sampler getSampler() const; - const property_list &getPropList() const; - -protected: - template - friend const decltype(Obj::impl) & - detail::getSyclObjImpl(const Obj &SyclObject); - - template - friend T detail::createSyclObjFromImpl( - std::add_rvalue_reference_t ImplObj); - - template - friend T detail::createSyclObjFromImpl( - std::add_lvalue_reference_t 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 class __image_array_slice__; diff --git a/sycl/include/sycl/accessor_image_base.hpp b/sycl/include/sycl/accessor_image_base.hpp new file mode 100644 index 0000000000000..69b22038e7495 --- /dev/null +++ b/sycl/include/sycl/accessor_image_base.hpp @@ -0,0 +1,172 @@ +//==--------------------- accessor_image_base.hpp --------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include +#include +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace detail { + +struct AccHostDataT; + +class UnsampledImageAccessorImplHost; +class SampledImageAccessorImplHost; +using UnsampledImageAccessorImplPtr = + std::shared_ptr; +using SampledImageAccessorImplPtr = + std::shared_ptr; + +class __SYCL_EXPORT UnsampledImageAccessorBaseHost { +protected: + UnsampledImageAccessorBaseHost(const UnsampledImageAccessorImplPtr &Impl) + : impl{Impl} {} + +public: + UnsampledImageAccessorBaseHost(sycl::range<3> Size, access_mode AccessMode, + void *SYCLMemObject, int Dims, int ElemSize, + id<3> Pitch, image_channel_type ChannelType, + image_channel_order ChannelOrder, + const property_list &PropertyList = {}); + const sycl::range<3> &getSize() const; + void *getMemoryObject() const; + detail::AccHostDataT &getAccData(); + void *getPtr(); + void *getPtr() const; + int getNumOfDims() const; + int getElementSize() const; + id<3> getPitch() const; + image_channel_type getChannelType() const; + image_channel_order getChannelOrder() const; + const property_list &getPropList() const; + +protected: + template + friend const decltype(Obj::impl) & + detail::getSyclObjImpl(const Obj &SyclObject); + + template + friend T detail::createSyclObjFromImpl( + std::add_rvalue_reference_t ImplObj); + + template + friend T detail::createSyclObjFromImpl( + std::add_lvalue_reference_t 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 { +protected: + SampledImageAccessorBaseHost(const SampledImageAccessorImplPtr &Impl) + : impl{Impl} {} + +public: + SampledImageAccessorBaseHost(sycl::range<3> Size, void *SYCLMemObject, + int Dims, int ElemSize, id<3> Pitch, + image_channel_type ChannelType, + image_channel_order ChannelOrder, + image_sampler Sampler, + const property_list &PropertyList = {}); + const sycl::range<3> &getSize() const; + void *getMemoryObject() const; + detail::AccHostDataT &getAccData(); + void *getPtr(); + void *getPtr() const; + int getNumOfDims() const; + int getElementSize() const; + id<3> getPitch() const; + image_channel_type getChannelType() const; + image_channel_order getChannelOrder() const; + image_sampler getSampler() const; + const property_list &getPropList() const; + +protected: + template + friend const decltype(Obj::impl) & + detail::getSyclObjImpl(const Obj &SyclObject); + + template + friend T detail::createSyclObjFromImpl( + std::add_rvalue_reference_t ImplObj); + + template + friend T detail::createSyclObjFromImpl( + std::add_lvalue_reference_t 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 +}; + +} // namespace detail +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/detail/accessor_impl.cpp b/sycl/source/detail/accessor_impl.cpp index 105d92e88660c..3c74f80638e93 100644 --- a/sycl/source/detail/accessor_impl.cpp +++ b/sycl/source/detail/accessor_impl.cpp @@ -39,11 +39,13 @@ void addHostAccessorAndWait(Requirement *Req) { Event->wait(); } -void addHostUnsampledImageAccessorAndWait(UnsampledImageAccessorImplHost *Req) { +void __SYCL_EXPORT +addHostUnsampledImageAccessorAndWait(UnsampledImageAccessorImplHost *Req) { addHostAccessorAndWait(Req); } -void addHostSampledImageAccessorAndWait(SampledImageAccessorImplHost *Req) { +void __SYCL_EXPORT +addHostSampledImageAccessorAndWait(SampledImageAccessorImplHost *Req) { addHostAccessorAndWait(Req); } @@ -55,7 +57,7 @@ void constructorNotification(void *BufferObj, void *AccessorObj, BufferObj, AccessorObj, (uint32_t)Target, (uint32_t)Mode, CodeLoc); } -void unsampledImageConstructorNotification( +void __SYCL_EXPORT unsampledImageConstructorNotification( void *ImageObj, void *AccessorObj, const std::optional &Target, access::mode Mode, const void *Type, uint32_t ElemSize, const code_location &CodeLoc) { @@ -68,7 +70,7 @@ void unsampledImageConstructorNotification( ImageObj, AccessorObj, (uint32_t)Mode, Type, ElemSize, CodeLoc); } -void sampledImageConstructorNotification( +void __SYCL_EXPORT sampledImageConstructorNotification( void *ImageObj, void *AccessorObj, const std::optional &Target, const void *Type, uint32_t ElemSize, const code_location &CodeLoc) { diff --git a/sycl/source/detail/accessor_impl.hpp b/sycl/source/detail/accessor_impl.hpp index 643247549e414..19899405d21c6 100644 --- a/sycl/source/detail/accessor_impl.hpp +++ b/sycl/source/detail/accessor_impl.hpp @@ -10,9 +10,9 @@ #include #include -#include #include #include +#include #include #include diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 71aa92124f1c8..5b8a5a917521f 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -27,6 +27,7 @@ #include #include #include +#include #include #include #include From 249cf63ad93813e3e48aa39044f150d171a02d81 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 15 Sep 2025 19:00:05 +0200 Subject: [PATCH 3/3] More gcc 7.5 fixes --- sycl/source/detail/image_accessor_util.cpp | 21 ++++++++++----------- 1 file changed, 10 insertions(+), 11 deletions(-) diff --git a/sycl/source/detail/image_accessor_util.cpp b/sycl/source/detail/image_accessor_util.cpp index 88ca0d1f58178..21670044308b3 100644 --- a/sycl/source/detail/image_accessor_util.cpp +++ b/sycl/source/detail/image_accessor_util.cpp @@ -6,8 +6,6 @@ // //===----------------------------------------------------------------------===// -#include -#include #include #include #include @@ -18,9 +16,9 @@ namespace detail { // For Nearest Filtering mode, process float4 Coordinates and // return the appropriate Pixel Coordinates based on Addressing Mode. -int4 getPixelCoordNearestFiltMode(float4 Coorduvw, - const addressing_mode SmplAddrMode, - const range<3> ImgRange) { +int4 __SYCL_EXPORT getPixelCoordNearestFiltMode( + float4 Coorduvw, const addressing_mode SmplAddrMode, + const range<3> ImgRange) { int4 Coordijk(0); int4 Rangewhd(ImgRange[0], ImgRange[1], ImgRange[2], 0); switch (SmplAddrMode) { @@ -91,9 +89,9 @@ int4 getPixelCoordNearestFiltMode(float4 Coorduvw, // Retabc contains the values of (a,b,c,0) // The caller of this function should use these values to create the 8 pixel // coordinates and multiplication coefficients. -int8 getPixelCoordLinearFiltMode(float4 Coorduvw, - const addressing_mode SmplAddrMode, - const range<3> ImgRange, float4 &Retabc) { +int8 __SYCL_EXPORT +getPixelCoordLinearFiltMode(float4 Coorduvw, const addressing_mode SmplAddrMode, + const range<3> ImgRange, float4 &Retabc) { int4 Rangewhd(ImgRange[0], ImgRange[1], ImgRange[2], 0); int4 Ci0j0k0(0); int4 Ci1j1k1(0); @@ -152,8 +150,9 @@ int8 getPixelCoordLinearFiltMode(float4 Coorduvw, // Note: For addressing_mode::none , spec says outofrange access is not defined. // This function handles this addressing_mode to avoid accessing out of bound // memories on host. -bool isOutOfRange(const int4 PixelCoord, const addressing_mode SmplAddrMode, - const range<3> ImgRange) { +bool __SYCL_EXPORT isOutOfRange(const int4 PixelCoord, + const addressing_mode SmplAddrMode, + const range<3> ImgRange) { if (SmplAddrMode != addressing_mode::clamp && SmplAddrMode != addressing_mode::none) @@ -170,7 +169,7 @@ bool isOutOfRange(const int4 PixelCoord, const addressing_mode SmplAddrMode, return (CheckWidth || CheckHeight || CheckDepth); } -float4 getBorderColor(const image_channel_order ImgChannelOrder) { +float4 __SYCL_EXPORT getBorderColor(const image_channel_order ImgChannelOrder) { float4 BorderColor(0.0f); switch (ImgChannelOrder) {