diff --git a/sycl/include/sycl/access/access.hpp b/sycl/include/sycl/access/access.hpp index b2c2eb3957b09..34c1d86a973ca 100644 --- a/sycl/include/sycl/access/access.hpp +++ b/sycl/include/sycl/access/access.hpp @@ -8,7 +8,7 @@ #pragma once -#include // for __SYCL2020_DEPRECATED +#include // for __SYCL2020_DEPRECATED, __SY... #ifdef __SYCL_DEVICE_ONLY__ #include @@ -106,8 +106,14 @@ template struct mode_target_tag_t { inline constexpr mode_tag_t read_only{}; inline constexpr mode_tag_t read_write{}; inline constexpr mode_tag_t write_only{}; + +// access::target::constant_buffer is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH +__SYCL2020_DEPRECATED("use 'read_only' instead") inline constexpr mode_target_tag_t read_constant{}; +__SYCL_SUPPRESS_DEPRECATED_POP + inline constexpr mode_target_tag_t read_only_host_task; inline constexpr mode_target_tag_t @@ -117,15 +123,6 @@ inline constexpr mode_target_tag_t namespace detail { -constexpr bool isTargetHostAccess(access::target T) { - return T == access::target::host_buffer || T == access::target::host_image; -} - -constexpr bool modeNeedsOldData(access::mode m) { - return m == access::mode::read || m == access::mode::write || - m == access::mode::read_write || m == access::mode::atomic; -} - constexpr bool modeWritesNewData(access::mode m) { return m != access::mode::read; } @@ -160,10 +157,14 @@ template <> struct TargetToAS { access::address_space::local_space; }; +// access::target::constant_buffer and address_space::constant_space are +// deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template <> struct TargetToAS { constexpr static access::address_space AS = access::address_space::constant_space; }; +__SYCL_SUPPRESS_DEPRECATED_POP template struct DecoratedType; @@ -183,6 +184,8 @@ struct DecoratedType { using type = __OPENCL_GLOBAL_AS__ ElementType; }; +// access::address_space::constant_space is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template struct DecoratedType { // Current implementation of address spaces handling leads to possibility @@ -196,6 +199,7 @@ struct DecoratedType { using type = __OPENCL_GLOBAL_AS__ ElementType; #endif }; +__SYCL_SUPPRESS_DEPRECATED_POP template struct DecoratedType { @@ -226,10 +230,13 @@ template struct deduce_AS_impl<__OPENCL_LOCAL_AS__ T> { access::address_space::local_space; }; +// access::address_space::constant_space is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template struct deduce_AS_impl<__OPENCL_CONSTANT_AS__ T> { static constexpr access::address_space value = access::address_space::constant_space; }; +__SYCL_SUPPRESS_DEPRECATED_POP template struct deduce_AS diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 0d8ea29488948..5b7bc3cc5755c 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -290,7 +290,11 @@ class accessor_common { protected: constexpr static access::address_space AS = TargetToAS::AS; + // access::target::host_buffer is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH constexpr static bool IsHostBuf = AccessTarget == access::target::host_buffer; + __SYCL_SUPPRESS_DEPRECATED_POP + constexpr static bool IsHostTask = AccessTarget == access::target::host_task; // SYCL2020 4.7.6.9.4.3 // IsPlaceHolder template parameter has no bearing on whether the accessor @@ -304,6 +308,9 @@ class accessor_common { // placeholder either. Otherwise, the accessor is a placeholder. constexpr static bool IsPlaceH = !IsHostBuf; + // access::target::global_buffer, constant_buffer and + // access::mode::discard_write, discard_read_write, atomic are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH // TODO: SYCL 2020 deprecates four of the target enum values // and replaces them with 2 (device and host_task). May want // to change these constexpr. @@ -326,6 +333,7 @@ class accessor_common { AccessMode == access::mode::read_write; constexpr static bool IsAccessAtomic = AccessMode == access::mode::atomic; + __SYCL_SUPPRESS_DEPRECATED_POP using RefType = detail::const_if_const_AS &; using ConstRefType = const DataT &; @@ -398,6 +406,8 @@ constexpr access::mode deduceAccessMode() { return access::mode::write; } + // access::target::constant_buffer is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH if constexpr (std::is_same_v< MayBeTag1, mode_target_tag_t>) { return access::mode::read; } + __SYCL_SUPPRESS_DEPRECATED_POP if constexpr (std::is_same_v constexpr access::target deduceAccessTarget(access::target defaultTarget) { + // access::target::constant_buffer is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH if constexpr (std::is_same_v< MayBeTag1, mode_target_tag_t>) { return access::target::constant_buffer; } + __SYCL_SUPPRESS_DEPRECATED_POP if constexpr ( std::is_same_v::value, "PropertyListT must be accessor_property_list"); @@ -679,6 +696,9 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : static access::mode getAdjustedMode(const PropertyListT &PropertyList) { access::mode AdjustedMode = AccessMode; + // property::noinit and access::mode::discard_write, discard_read_write + // are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH if (PropertyList.template has_property() || PropertyList.template has_property()) { if (AdjustedMode == access::mode::write) { @@ -687,6 +707,7 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : AdjustedMode = access::mode::discard_read_write; } } + __SYCL_SUPPRESS_DEPRECATED_POP return AdjustedMode; } @@ -1666,6 +1687,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : return *(getQualifiedPtr() + LinearIndex); } + // access_mode::atomic is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template > @@ -1681,6 +1704,7 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : *getQualifiedPtr() = std::move(Other); return *this; } + __SYCL_SUPPRESS_DEPRECATED_POP template 0) && @@ -1690,6 +1714,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : return getQualifiedPtr()[LinearIndex]; } + // access::mode::atomic and sycl::atomic are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template operator typename std::enable_if_t< Dims == 0 && AccessMode == access::mode::atomic, atomic>() @@ -1716,11 +1742,16 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : return atomic(multi_ptr( getQualifiedPtr() + LinearIndex)); } + __SYCL_SUPPRESS_DEPRECATED_POP + template 1)>> auto operator[](size_t Index) const { return AccessorSubscript(*this, Index); } + // access::target::host_buffer, constant_buffer and constant_ptr are + // deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template get_pointer() const { return constant_ptr(getPointerAdjusted()); } + __SYCL_SUPPRESS_DEPRECATED_POP template typename std::enable_if_t< @@ -2052,6 +2085,8 @@ accessor(buffer, handler &, Type1, Type2, Type3, access::placeholder::false_t, ext::oneapi::accessor_property_list>; +// access::target::local is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH /// Local accessor /// /// \ingroup sycl_api_acc @@ -2069,6 +2104,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : using AccessorCommonT = detail::accessor_common; + __SYCL_SUPPRESS_DEPRECATED_POP using AccessorCommonT::AS; @@ -2171,6 +2207,8 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : using reference = DataT &; using const_reference = const DataT &; + // access::target::local is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template > local_accessor_base(handler &, const detail::code_location CodeLoc = detail::code_location::current()) @@ -2241,6 +2279,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : } #endif + __SYCL_SUPPRESS_DEPRECATED_POP size_t get_size() const { return getSize().size() * sizeof(DataT); } __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") @@ -2274,6 +2313,8 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : return getQualifiedPtr()[Index]; } + // access::mode::atomic and sycl::atomic are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template operator typename std::enable_if_t< Dims == 0 && AccessMode == access::mode::atomic, atomic>() @@ -2298,6 +2339,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : return atomic(multi_ptr( getQualifiedPtr() + Index)); } + __SYCL_SUPPRESS_DEPRECATED_POP template 1)>> typename AccessorCommonT::template AccessorSubscript< @@ -2316,6 +2358,8 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : }; // TODO: Remove deprecated specialization once no longer needed +// access::target::local is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template class __SYCL_EBO __SYCL_SPECIAL_CLASS accessor< @@ -2325,6 +2369,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS accessor< accessor> { + __SYCL_SUPPRESS_DEPRECATED_POP + using local_acc = local_accessor_base; @@ -2547,6 +2593,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor friend class ext::oneapi::experimental::dynamic_local_accessor; }; +// access::target::host_buffer is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template class __SYCL_EBO host_accessor @@ -2573,12 +2621,12 @@ class __SYCL_EBO host_accessor id Offset) { AccessorT::__init(Ptr, AccessRange, MemRange, Offset); } - #ifndef __SYCL_DEVICE_ONLY__ host_accessor(const detail::AccessorImplPtr &Impl) : accessor{Impl} {} #endif // __SYCL_DEVICE_ONLY__ + __SYCL_SUPPRESS_DEPRECATED_POP public: host_accessor() : AccessorT() {} @@ -2721,6 +2769,8 @@ class __SYCL_EBO host_accessor : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset, PropertyList, CodeLoc) {} + // access_mode::atomic is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template > @@ -2737,6 +2787,7 @@ class __SYCL_EBO host_accessor *AccessorT::getQualifiedPtr() = std::move(Other); return *this; } + __SYCL_SUPPRESS_DEPRECATED_POP // implicit conversion between const / non-const types for read only accessors template +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP #include #include #include @@ -220,6 +221,10 @@ class image_accessor template friend class __image_array_slice__; + // access::target::host_image, image, image_array and mode::discard_write are + // deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH + constexpr static bool IsHostImageAcc = (AccessTarget == access::target::host_image); @@ -232,6 +237,8 @@ class image_accessor (AccessMode == access::mode::write || AccessMode == access::mode::discard_write); + __SYCL_SUPPRESS_DEPRECATED_POP + constexpr static bool IsImageAccessAnyWrite = (IsImageAccessWriteOnly || AccessMode == access::mode::read_write); @@ -433,6 +440,9 @@ class image_accessor #endif + // sycl::sampler is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH + // Available only when: // (accessTarget == access::target::image && accessMode == access::mode::read) // || (accessTarget == access::target::host_image && ( accessMode == @@ -475,6 +485,8 @@ class image_accessor #endif } + __SYCL_SUPPRESS_DEPRECATED_POP + // Available only when: // (accessTarget == access::target::image && (accessMode == // access::mode::write || accessMode == access::mode::discard_write)) || @@ -528,6 +540,9 @@ class __image_array_slice__ { } public: + // access::target::image_array and sycl::sampler are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH + __image_array_slice__( accessor> @@ -548,6 +563,8 @@ class __image_array_slice__ { return MBaseAcc.read(getAdjustedCoords(Coords), Smpl); } + __SYCL_SUPPRESS_DEPRECATED_POP + template ::value>> void write(const CoordT &Coords, const DataT &Color) const { @@ -588,12 +605,20 @@ class __image_array_slice__ { private: size_t MIdx; + + // access::target::image_array is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH + accessor> MBaseAcc; + __SYCL_SUPPRESS_DEPRECATED_POP }; } // namespace detail + +// access::target::image and sycl::image are deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH /// Image accessors. /// /// Available only when accessTarget == access::target::image. @@ -656,7 +681,10 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor< accessor() = default; #endif }; +__SYCL_SUPPRESS_DEPRECATED_POP +// access::target::host_image and sycl::image are deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH /// Host image accessor. /// /// Available only when accessTarget == access::target::host_image. @@ -689,7 +717,10 @@ class __SYCL_EBO accessor // for __SYCL2020_DEPRECATED +#include // for __SYCL2020_DEPRECATED, __SY... #include // for uint8_t, int16_t, int32_t @@ -137,11 +137,14 @@ using cl_float = float; using cl_double = double; } // namespace opencl +// sycl::cl_* scalar aliases used by the vector alias macros are deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH __SYCL_2020_MAKE_VECTOR_ALIASES_FOR_VECTOR_LENGTH(2) __SYCL_2020_MAKE_VECTOR_ALIASES_FOR_VECTOR_LENGTH(3) __SYCL_2020_MAKE_VECTOR_ALIASES_FOR_VECTOR_LENGTH(4) __SYCL_2020_MAKE_VECTOR_ALIASES_FOR_VECTOR_LENGTH(8) __SYCL_2020_MAKE_VECTOR_ALIASES_FOR_VECTOR_LENGTH(16) +__SYCL_SUPPRESS_DEPRECATED_POP } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/backend_types.hpp b/sycl/include/sycl/backend_types.hpp index e80a01bcf1d78..446a2b00c4b14 100644 --- a/sycl/include/sycl/backend_types.hpp +++ b/sycl/include/sycl/backend_types.hpp @@ -8,7 +8,7 @@ #pragma once -#include // for __SYCL2020_DEPRECATED +#include // for __SYCL2020_DEPRECATED, __SY... #include // for __SYCL_EXPORT #include // for ostream @@ -44,8 +44,11 @@ __SYCL_EXPORT std::ostream &operator<<(std::ostream &Out, backend be); namespace detail { inline std::string_view get_backend_name_no_vendor(backend Backend) { switch (Backend) { + // backend::host is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH case backend::host: return "host"; + __SYCL_SUPPRESS_DEPRECATED_POP case backend::opencl: return "opencl"; case backend::ext_oneapi_level_zero: diff --git a/sycl/include/sycl/detail/address_space_cast.hpp b/sycl/include/sycl/detail/address_space_cast.hpp index 95a48e6474cd9..4a6270cc72853 100644 --- a/sycl/include/sycl/detail/address_space_cast.hpp +++ b/sycl/include/sycl/detail/address_space_cast.hpp @@ -10,6 +10,7 @@ #include #include +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP #include @@ -22,7 +23,11 @@ inline constexpr bool address_space_cast_is_possible(access::address_space Src, access::address_space Dst) { // constant_space is unique and is not interchangeable with any other. + // access::address_space::constant_space is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH auto constant_space = access::address_space::constant_space; + __SYCL_SUPPRESS_DEPRECATED_POP + if (Src == constant_space || Dst == constant_space) return Src == Dst; diff --git a/sycl/include/sycl/detail/helpers.hpp b/sycl/include/sycl/detail/helpers.hpp index 9110b32584a00..b34ae5f2b4f1e 100644 --- a/sycl/include/sycl/detail/helpers.hpp +++ b/sycl/include/sycl/detail/helpers.hpp @@ -9,9 +9,10 @@ #pragma once #include // for MemorySemanticsMask -#include // for fence_space -#include // for __SYCL_EXPORT -#include // for memory_order +#include // for fence_space +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP +#include // for __SYCL_EXPORT +#include // for memory_order #ifdef __SYCL_DEVICE_ONLY__ #include @@ -186,6 +187,8 @@ getSPIRVMemorySemanticsMask(memory_order) { } #ifndef __INTEL_PREVIEW_BREAKING_CHANGES +// access::fence_space is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH inline constexpr uint32_t getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace, const __spv::MemorySemanticsMask LocalScopeMask = @@ -230,6 +233,7 @@ getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace, __spv::MemorySemanticsMask::CrossWorkgroupMemory | LocalScopeMask); } +__SYCL_SUPPRESS_DEPRECATED_POP #endif // __INTEL_PREVIEW_BREAKING_CHANGES inline constexpr bool is_power_of_two(int x) { return (x & (x - 1)) == 0; } diff --git a/sycl/include/sycl/detail/image_ocl_types.hpp b/sycl/include/sycl/detail/image_ocl_types.hpp index 04329bea7c96b..35367c77585e1 100644 --- a/sycl/include/sycl/detail/image_ocl_types.hpp +++ b/sycl/include/sycl/detail/image_ocl_types.hpp @@ -32,6 +32,7 @@ #include #include +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP namespace sycl { inline namespace _V1 { @@ -552,6 +553,8 @@ inline int getSPIRVElementSize(int ImageChannelType, int ImageChannelOrder) { template struct opencl_image_type; +// access::target::host_image is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH // Creation of dummy ocl types for host_image targets. // These dummy ocl types are needed by the compiler parser for the compilation // of host application code with __SYCL_DEVICE_ONLY__ macro set. @@ -560,6 +563,8 @@ struct opencl_image_type { using type = opencl_image_type *; }; +__SYCL_SUPPRESS_DEPRECATED_POP + template struct sampled_opencl_image_type { using type = void *; }; diff --git a/sycl/include/sycl/ext/oneapi/experimental/fragment.hpp b/sycl/include/sycl/ext/oneapi/experimental/fragment.hpp index ae6b4b531960d..39256b85bf415 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/fragment.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/fragment.hpp @@ -185,7 +185,9 @@ __SYCL_ALWAYS_INLINE inline fragment get_opportunistic_group() { #ifdef __SYCL_DEVICE_ONLY__ #if defined(__SPIR__) || defined(__SPIRV__) - sycl::sub_group sg = sycl::ext::oneapi::experimental::this_sub_group(); + + sycl::sub_group sg = sycl::ext::oneapi::this_work_item::get_sub_group(); + sub_group_mask mask = sycl::ext::oneapi::group_ballot(sg, true); return fragment(mask, 0, 1); #elif defined(__NVPTX__) diff --git a/sycl/include/sycl/ext/oneapi/filter_selector.hpp b/sycl/include/sycl/ext/oneapi/filter_selector.hpp index ab5b33324fd30..52913b6b70780 100644 --- a/sycl/include/sycl/ext/oneapi/filter_selector.hpp +++ b/sycl/include/sycl/ext/oneapi/filter_selector.hpp @@ -8,9 +8,10 @@ #pragma once -#include // for __SYCL_EXPORT -#include // for device -#include // for device_selector +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP +#include // for __SYCL_EXPORT +#include // for device +#include // for device_selector #include // for shared_ptr #include // for string @@ -34,6 +35,8 @@ namespace detail { class filter_selector_impl; } // namespace detail +// sycl::device_selector and ext::oneapi::filter_selector are deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH class __SYCL_EXPORT_DEPRECATED( "filter_selector is deprecated, use a SYCL 2020 device selector callable " "instead") filter_selector : public device_selector { @@ -51,6 +54,9 @@ class __SYCL_EXPORT_DEPRECATED( std::shared_ptr impl; filter_selector(sycl::detail::string_view filter); }; + +__SYCL_SUPPRESS_DEPRECATED_POP + } // namespace ext::oneapi #ifdef __SYCL_INTERNAL_API diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp index c5f93cd5fe957..b28d836b96da1 100644 --- a/sycl/include/sycl/group.hpp +++ b/sycl/include/sycl/group.hpp @@ -13,7 +13,7 @@ #include #include // for convertToOpenCLGroupAsyncCopyPtr #include // for __SYCL_TYPE -#include // for __SYCL2020_DEPRECATED +#include // for __SYCL2020_DEPRECATED, __SY... #include #include // for Builder, getSPIRVMemo... #include // for NDLoop @@ -307,6 +307,8 @@ template class __SYCL_TYPE(group) group { } #ifndef __INTEL_PREVIEW_BREAKING_CHANGES + // access::fence_space is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH /// Executes a work-group mem-fence with memory ordering on the local address /// space, global address space or both based on the value of \p accessSpace. template @@ -331,6 +333,7 @@ template class __SYCL_TYPE(group) group { __spirv_MemoryBarrier(__spv::Scope::Workgroup, flags); #endif } + __SYCL_SUPPRESS_DEPRECATED_POP #endif // __INTEL_PREVIEW_BREAKING_CHANGES /// Asynchronously copies a number of elements specified by \p numElements diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 1be6bce3d1e4e..97e317e48b103 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -12,7 +12,7 @@ #include #include #include -#include +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP #include #include #include @@ -473,6 +473,8 @@ class __SYCL_EXPORT handler { void setArgsHelper(int) {} + // access::target::local is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH void setLocalAccessorArgHelper(int ArgIndex, detail::LocalAccessorBaseHost &LocalAccBase) { detail::LocalAccessorImplPtr LocalAccImpl = @@ -495,6 +497,7 @@ class __SYCL_EXPORT handler { setLocalAccessorArgHelper(ArgIndex, Arg); #endif } + __SYCL_SUPPRESS_DEPRECATED_POP // setArgHelper for local accessor argument (up to date accessor interface) template @@ -509,6 +512,8 @@ class __SYCL_EXPORT handler { void setArgHelper(int ArgIndex, detail::work_group_memory_impl &Arg); // setArgHelper for non local accessor argument. + // access::target::local is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template typename std::enable_if_t @@ -524,6 +529,7 @@ class __SYCL_EXPORT handler { addArg(detail::kernel_param_kind_t::kind_accessor, Req, static_cast(AccessTarget), ArgIndex); } + __SYCL_SUPPRESS_DEPRECATED_POP template void setArgHelper(int ArgIndex, T &&Arg) { void *StoredArg = storePlainArg(Arg); @@ -541,11 +547,14 @@ class __SYCL_EXPORT handler { } } + // sycl::sampler is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH void setArgHelper(int ArgIndex, sampler &&Arg) { void *StoredArg = storePlainArg(Arg); addArg(detail::kernel_param_kind_t::kind_sampler, StoredArg, sizeof(sampler), ArgIndex); } + __SYCL_SUPPRESS_DEPRECATED_POP void setArgHelper(int ArgIndex, stream &&Str); @@ -764,6 +773,8 @@ class __SYCL_EXPORT handler { return false; } + // access::target::constant_buffer, image and image_array are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH constexpr static bool isConstOrGlobal(access::target AccessTarget) { return AccessTarget == access::target::device || AccessTarget == access::target::constant_buffer; @@ -773,6 +784,7 @@ class __SYCL_EXPORT handler { return AccessTarget == access::target::image || AccessTarget == access::target::image_array; } + __SYCL_SUPPRESS_DEPRECATED_POP constexpr static bool isValidTargetForExplicitOp(access::target AccessTarget) { @@ -784,6 +796,8 @@ class __SYCL_EXPORT handler { AccessMode == access::mode::read_write; } + // access::mode::discard_write and discard_read_write are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH constexpr static bool isValidModeForDestinationAccessor(access::mode AccessMode) { return AccessMode == access::mode::write || @@ -791,6 +805,7 @@ class __SYCL_EXPORT handler { AccessMode == access::mode::discard_write || AccessMode == access::mode::discard_read_write; } + __SYCL_SUPPRESS_DEPRECATED_POP // UR APIs only support select fill sizes: 1, 2, 4, 8, 16, 32, 64, 128 constexpr static bool isBackendSupportedFillSize(size_t Size) { @@ -1166,6 +1181,8 @@ class __SYCL_EXPORT handler { template using is_same_type = std::is_same, remove_cv_ref_t>; + // sycl::sampler is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template struct ShouldEnableSetArg { static constexpr bool value = std::is_trivially_copyable_v> @@ -1179,6 +1196,7 @@ class __SYCL_EXPORT handler { || is_same_type::value // Stream || detail::check_if_device_copyable_v>; }; + __SYCL_SUPPRESS_DEPRECATED_POP /// Sets argument for OpenCL interoperability kernels. /// @@ -1214,6 +1232,9 @@ class __SYCL_EXPORT handler { constexpr int AccessTargetMask = 0x7ff; const access::target target = static_cast( type::sizes[NumArgs] & AccessTargetMask); + + // access::target::local is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH if (target == target::local) { detail::LocalAccessorBaseHost *LocalAccBase = (detail::LocalAccessorBaseHost *)(FieldArg); @@ -1227,6 +1248,7 @@ class __SYCL_EXPORT handler { addArg(type::kinds[NumArgs], Req, type::sizes[NumArgs], ArgIndex + NumArgs); } + __SYCL_SUPPRESS_DEPRECATED_POP } else { // for non-accessors, simply call addArg normally. addArg(type::kinds[NumArgs], FieldArg, type::sizes[NumArgs], @@ -2925,6 +2947,8 @@ class __SYCL_EXPORT handler { const detail::kernel_param_kind_t &Kind = ParamDesc.kind; const access::target AccTarget = static_cast(ParamDesc.info & AccessTargetMask); + // access::target::local is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH if ((Kind == detail::kernel_param_kind_t::kind_accessor) && (AccTarget == target::local)) throw sycl::exception( @@ -2932,6 +2956,7 @@ class __SYCL_EXPORT handler { "A local accessor must not be used in a SYCL kernel function " "that is invoked via single_task or via the simple form of " "parallel_for that takes a range parameter."); + __SYCL_SUPPRESS_DEPRECATED_POP if (Kind == detail::kernel_param_kind_t::kind_work_group_memory || Kind == detail::kernel_param_kind_t::kind_dynamic_work_group_memory) throw sycl::exception( diff --git a/sycl/include/sycl/image.hpp b/sycl/include/sycl/image.hpp index a43db0f08c4be..dbe627519947a 100644 --- a/sycl/include/sycl/image.hpp +++ b/sycl/include/sycl/image.hpp @@ -107,7 +107,10 @@ enum class image_format : unsigned int { using byte = unsigned char; +// sycl::byte is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH using image_allocator = detail::aligned_allocator; +__SYCL_SUPPRESS_DEPRECATED_POP namespace detail { @@ -759,6 +762,8 @@ class unsampled_image #endif public: + // sycl::byte is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH unsampled_image( image_format Format, const range &Range, const property_list &PropList = {}, @@ -967,6 +972,7 @@ class unsampled_image CodeLoc, (void *)this->impl.get(), HostPointer.get(), Dimensions, detail::rangeToArray(Range).data(), Format); } + __SYCL_SUPPRESS_DEPRECATED_POP #ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS template &Range, const property_list &PropList = {}, @@ -1241,6 +1249,7 @@ class sampled_image CodeLoc, (void *)this->impl.get(), HostPointer.get(), Dimensions, detail::rangeToArray(Range).data(), Format, Sampler); } + __SYCL_SUPPRESS_DEPRECATED_POP #ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS template // for target, mode, place... -#include // for AccessorBaseHost -#include // for backend, backend_re... -#include // for buffer -#include // for __SYCL_EXPORT -#include // for getSyclObjImpl +#include // for target, mode, place... +#include // for AccessorBaseHost +#include // for backend, backend_re... +#include // for buffer +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP +#include // for __SYCL_EXPORT +#include // for getSyclObjImpl #include #include // for accessor_property_list #include // for command_graph @@ -54,6 +55,8 @@ class interop_handle { /// a backend graph object is available for interop. __SYCL_EXPORT bool ext_codeplay_has_graph() const noexcept; + // access::target::constant_buffer is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH /// Receives a SYCL accessor that has been defined as a requirement for the /// command group, and returns the underlying OpenCL memory object that is /// used by the SYCL runtime. If the accessor passed as parameter is not part @@ -84,6 +87,8 @@ class interop_handle { #endif } + __SYCL_SUPPRESS_DEPRECATED_POP + /// Receives a SYCL accessor that has been defined as a requirement for the /// command group, and returns the underlying OpenCL memory object that is /// used by the SYCL runtime. If the accessor passed as parameter is not part diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp index 97ea1d7b3e864..8fbfbd81e4506 100644 --- a/sycl/include/sycl/multi_ptr.hpp +++ b/sycl/include/sycl/multi_ptr.hpp @@ -10,7 +10,7 @@ #include #include -#include +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP #include #include #include @@ -37,13 +37,18 @@ struct LegacyPointerTypes { // Specialization for constant_space to avoid creating a non-legacy multi_ptr // with the unsupported space. +// access::address_space::constant_space is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template struct LegacyPointerTypes { + using decorated_type = typename detail::DecoratedType< ElementType, access::address_space::constant_space>::type; + using pointer_t = decorated_type *; using const_pointer_t = decorated_type const *; }; +__SYCL_SUPPRESS_DEPRECATED_POP // Helper to avoid instantiations of invalid non-legacy multi_ptr types. template @@ -57,14 +62,19 @@ struct LegacyReferenceTypes { // Specialization for constant_space to avoid creating a non-legacy multi_ptr // with the unsupported space. +// access::address_space::constant_space is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template struct LegacyReferenceTypes { + using decorated_type = typename detail::DecoratedType< ElementType, access::address_space::constant_space>::type; + using reference_t = decorated_type &; using const_reference_t = decorated_type &; }; +__SYCL_SUPPRESS_DEPRECATED_POP } // namespace detail /// Provides constructors for address space qualified and non address space @@ -258,6 +268,9 @@ class __SYCL_TYPE(multi_ptr) multi_ptr { m_Pointer = nullptr; return *this; } + + // access::address_space::constant_space is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template < access::address_space OtherSpace, access::decorated OtherIsDecorated, typename = @@ -268,6 +281,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr { m_Pointer = detail::static_address_cast(Other.get_decorated()); return *this; } + template < access::address_space OtherSpace, access::decorated OtherIsDecorated, typename = @@ -278,6 +292,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr { m_Pointer = detail::static_address_cast(std::move(Other.m_Pointer)); return *this; } + __SYCL_SUPPRESS_DEPRECATED_POP reference operator*() const { return *m_Pointer; } pointer operator->() const { return get(); } @@ -562,10 +577,14 @@ class __SYCL_TYPE(multi_ptr) multi_ptr { std::add_pointer_t>); // Legacy has a different interface. static_assert(DecorateAddress != access::decorated::legacy); + // constant_space is only supported in legacy multi_ptr. + // access::address_space::constant_space is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH static_assert(Space != access::address_space::constant_space, "SYCL 2020 multi_ptr does not support the deprecated " "constant_space address space."); + __SYCL_SUPPRESS_DEPRECATED_POP // Constructors multi_ptr() : m_Pointer(nullptr) {} @@ -813,6 +832,9 @@ class multi_ptr { Accessor.get_pointer().get())) {} // Only if Space == local_space || generic_space + // target::local, constant_buffer, address_space::constant_space and + // accessor::get_pointer() are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template < int dimensions, access::mode Mode, access::placeholder isPlaceholder, typename PropertyListT, access::address_space _Space = Space, @@ -904,6 +926,7 @@ class multi_ptr { target::constant_buffer, isPlaceholder, PropertyListT> Accessor) : multi_ptr(Accessor.get_pointer()) {} + __SYCL_SUPPRESS_DEPRECATED_POP // TODO: This constructor is the temporary solution for the existing problem // with conversions from multi_ptr { : multi_ptr(Accessor.get_pointer()) {} // Only if Space == local_space || generic_space + // target::local, constant_buffer and address_space::constant_space are + // deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template < typename ElementType, int dimensions, access::mode Mode, typename PropertyListT, access::address_space _Space = Space, @@ -1141,6 +1167,7 @@ class multi_ptr { access::placeholder::false_t, PropertyListT> Accessor) : multi_ptr(Accessor.get_pointer()) {} + __SYCL_SUPPRESS_DEPRECATED_POP using ReturnPtr = detail::const_if_const_AS *; // Returns the underlying OpenCL C pointer @@ -1269,6 +1296,9 @@ class multi_ptr { : multi_ptr(Accessor.get_pointer()) {} // Only if Space == local_space || generic_space + // target::local, constant_buffer and address_space::constant_space are + // deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH template < typename ElementType, int dimensions, access::mode Mode, typename PropertyListT, access::address_space _Space = Space, @@ -1300,6 +1330,7 @@ class multi_ptr { access::placeholder::false_t, PropertyListT> Accessor) : multi_ptr(Accessor.get_pointer()) {} + __SYCL_SUPPRESS_DEPRECATED_POP // Returns the underlying OpenCL C pointer pointer_t get() const { return m_Pointer; } @@ -1348,17 +1379,24 @@ template ) -> multi_ptr; + +// target::constant_buffer, local and address_space::constant_space are +// deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template multi_ptr(accessor) -> multi_ptr; + template multi_ptr( accessor) -> multi_ptr; +__SYCL_SUPPRESS_DEPRECATED_POP + template multi_ptr(local_accessor) -> multi_ptr; diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp index 2adc9ecd6b122..bc54ee80764ae 100644 --- a/sycl/include/sycl/nd_item.hpp +++ b/sycl/include/sycl/nd_item.hpp @@ -71,7 +71,12 @@ template class nd_item { size_t LinId = 0; id Index = get_global_id(); range Extent = get_global_range(); + + // nd_item::get_offset() is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH id Offset = get_offset(); + __SYCL_SUPPRESS_DEPRECATED_POP + if (1 == Dimensions) { LinId = Index[0] - Offset[0]; } else if (2 == Dimensions) { @@ -198,11 +203,18 @@ template class nd_item { } nd_range get_nd_range() const noexcept { + // nd_item::get_offset() and the offset-taking nd_range constructor + // are deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH return nd_range(get_global_range(), get_local_range(), get_offset()); + __SYCL_SUPPRESS_DEPRECATED_POP } #ifndef __INTEL_PREVIEW_BREAKING_CHANGES + + // access::fence_space is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH __SYCL2020_DEPRECATED("use sycl::group_barrier() free function instead") void barrier([[maybe_unused]] access::fence_space accessSpace = access::fence_space::global_and_local) const { @@ -237,6 +249,7 @@ template class nd_item { __spirv_MemoryBarrier(__spv::Scope::Workgroup, flags); #endif } + __SYCL_SUPPRESS_DEPRECATED_POP #endif // __INTEL_PREVIEW_BREAKING_CHANGES /// Asynchronously copies a number of elements specified by \p numElements diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp index 6352ad79a4ff8..d1a2734461ee2 100644 --- a/sycl/include/sycl/reduction.hpp +++ b/sycl/include/sycl/reduction.hpp @@ -8,26 +8,27 @@ #pragma once -#include // for address_s... -#include // for local_acc... -#include // for aspect -#include // for IsValidAt... -#include // for atomic_ref -#include // for buffer -#include // for min -#include // for __SYCL_EX... -#include // for is_sgenfloat -#include // for createSyc... -#include // for id -#include // for strategy -#include // for make_tuple -#include // for device -#include // for event -#include // for make_erro... -#include // for queue_impl -#include // for plus, known_op... -#include // for workGroup... -#include // for reduce_ov... +#include // for address_s... +#include // for local_acc... +#include // for aspect +#include // for IsValidAt... +#include // for atomic_ref +#include // for buffer +#include // for min +#include // for __SYCL_SU... +#include // for __SYCL_EX... +#include // for is_sgenfloat +#include // for createSyc... +#include // for id +#include // for strategy +#include // for make_tuple +#include // for device +#include // for event +#include // for make_erro... +#include // for queue_impl +#include // for plus, known_op... +#include // for workGroup... +#include // for reduce_ov... #include #include // for handler #include // for getDeline... @@ -1357,11 +1358,15 @@ struct NDRangeReduction< // Integrated/discrete GPUs have different faster path. For discrete GPUs // fast path requires USM device allocations though, so check for that as // well. + // info::device::host_unified_memory is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH if (device.get_info() || !device.has(aspect::usm_device_allocations)) Rest(Redu.getReadWriteAccessorToInitializedGroupsCounter(CGH)); else Rest(Redu.getGroupsCounterAccDiscrete(CGH)); + + __SYCL_SUPPRESS_DEPRECATED_POP } }; diff --git a/sycl/include/sycl/sampler.hpp b/sycl/include/sycl/sampler.hpp index 9d77a6eb7501b..10ded14e0ef16 100644 --- a/sycl/include/sycl/sampler.hpp +++ b/sycl/include/sycl/sampler.hpp @@ -10,7 +10,8 @@ #include // for mode, placeholder, target #include // for __SYCL_SPECIAL_CLASS, __SYCL_TYPE -#include // for __SYCL_EXPORT +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP +#include // for __SYCL_EXPORT #include #include // for getSyclObjImpl #include // for property_list @@ -141,7 +142,10 @@ struct image_sampler { } // namespace _V1 } // namespace sycl +// sycl::sampler is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH template <> struct std::hash : public sycl::detail::sycl_obj_hash {}; +__SYCL_SUPPRESS_DEPRECATED_POP diff --git a/sycl/include/sycl/stream.hpp b/sycl/include/sycl/stream.hpp index 5e914af793e42..a184c4613d4aa 100644 --- a/sycl/include/sycl/stream.hpp +++ b/sycl/include/sycl/stream.hpp @@ -15,7 +15,7 @@ #include // for isinf, isnan, signbit #include // for array #include // for __SYCL_SPECIAL_CLASS, __S... -#include // for __SYCL2020_DEPRECATED +#include // for __SYCL2020_DEPRECATED, __SY... #include // for __SYCL_EXPORT #include // for OwnerLessBase #include @@ -97,9 +97,13 @@ using GlobalBufPtrType = typename detail::DecoratedType::type *; constexpr static int GlobalBufDim = 1; +// access::mode::atomic is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH using GlobalOffsetAccessorT = accessor; +__SYCL_SUPPRESS_DEPRECATED_POP + constexpr static access::address_space GlobalOffsetAS = TargetToAS::AS; using GlobalOffsetPtrType = @@ -644,6 +648,8 @@ writeFloatingPoint(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, : 0); } +// sycl::atomic, returned by the atomic accessor, is deprecated. +__SYCL_SUPPRESS_DEPRECATED_PUSH // Helper method to update offset in the global buffer atomically according to // the provided size of the data in the flush buffer. Return true if offset is // updated and false in case of overflow. @@ -661,6 +667,8 @@ inline bool updateOffset(GlobalOffsetAccessorT &GlobalOffset, return true; } +__SYCL_SUPPRESS_DEPRECATED_POP + inline void flushBuffer(GlobalOffsetAccessorT &GlobalOffset, GlobalBufAccessorT &GlobalBuf, GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset) { diff --git a/sycl/include/sycl/sycl_span.hpp b/sycl/include/sycl/sycl_span.hpp index d7f8a5971a591..32fffc57ef2eb 100644 --- a/sycl/include/sycl/sycl_span.hpp +++ b/sycl/include/sycl/sycl_span.hpp @@ -142,12 +142,13 @@ template */ -#include // for array -#include // for assert -#include // for size_t, nullptr_t, ptrdiff_t -#include // for SIZE_MAX -#include // for size, data, reverse_iterator -#include // for pointer_traits. +#include // for array +#include // for assert +#include // for size_t, nullptr_t, ptrdiff_t +#include // for SIZE_MAX +#include // for size, data, reverse_iterator +#include // for pointer_traits. +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP #include // for enable_if_t, enable_if, remove_cv_t, false_type #include // for declval @@ -480,6 +481,8 @@ template class _SYCL_SPAN_TEMPLATE_VIS span { return reverse_iterator(begin()); } + // sycl::byte is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH _SYCL_SPAN_INLINE_VISIBILITY span __as_bytes() const noexcept { return span{ @@ -491,6 +494,7 @@ template class _SYCL_SPAN_TEMPLATE_VIS span { return span{ reinterpret_cast(data()), size_bytes()}; } + __SYCL_SUPPRESS_DEPRECATED_POP private: pointer __data; @@ -694,6 +698,8 @@ class _SYCL_SPAN_TEMPLATE_VIS span<_Tp, dynamic_extent> { return reverse_iterator(begin()); } + // sycl::byte is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH _SYCL_SPAN_INLINE_VISIBILITY span __as_bytes() const noexcept { return {reinterpret_cast(data()), size_bytes()}; @@ -703,6 +709,7 @@ class _SYCL_SPAN_TEMPLATE_VIS span<_Tp, dynamic_extent> { __as_writable_bytes() const noexcept { return {reinterpret_cast(data()), size_bytes()}; } + __SYCL_SUPPRESS_DEPRECATED_POP private: pointer __data; diff --git a/sycl/test-e2e/Basic/accessor/Inputs/device_accessor.cpp b/sycl/test-e2e/Basic/accessor/Inputs/device_accessor.cpp index ca51c57755cbd..530fbab0b1ad9 100644 --- a/sycl/test-e2e/Basic/accessor/Inputs/device_accessor.cpp +++ b/sycl/test-e2e/Basic/accessor/Inputs/device_accessor.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include #include +#include // for __SYCL_SUPPRESS_DEPRECATED_PUSH/POP int main() { // Non-placeholder accessors. @@ -171,11 +172,14 @@ int main() { sycl::accessor acc_2(buf_data, cgh, sycl::range<1>(8), sycl::no_init); sycl::accessor acc_3(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1), sycl::no_init); + // sycl::read_constant is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH sycl::accessor acc_4(buf_data, cgh, sycl::read_constant); sycl::accessor acc_5(buf_data, cgh, sycl::range<1>(8), sycl::read_constant); sycl::accessor acc_6(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1), sycl::read_constant); + __SYCL_SUPPRESS_DEPRECATED_POP sycl::accessor acc_7(buf_data, cgh, sycl::write_only, sycl::no_init); sycl::accessor acc_8(buf_data, cgh, sycl::range<1>(8), sycl::write_only, sycl::no_init); @@ -186,11 +190,14 @@ int main() { auto acc_2 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::no_init); auto acc_3 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1), sycl::no_init); + // sycl::read_constant is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH auto acc_4 = buf_data.get_access(cgh, sycl::read_constant); auto acc_5 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::read_constant); auto acc_6 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1), sycl::read_constant); + __SYCL_SUPPRESS_DEPRECATED_POP auto acc_7 = buf_data.get_access(cgh, sycl::write_only, sycl::no_init); auto acc_8 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::write_only, sycl::no_init); @@ -241,10 +248,13 @@ int main() { sycl::accessor acc_2(buf_data, sycl::range<1>(8), sycl::no_init); sycl::accessor acc_3(buf_data, sycl::range<1>(8), sycl::id<1>(1), sycl::no_init); + // sycl::read_constant is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH sycl::accessor acc_4(buf_data, sycl::read_constant); sycl::accessor acc_5(buf_data, sycl::range<1>(8), sycl::read_constant); sycl::accessor acc_6(buf_data, sycl::range<1>(8), sycl::id<1>(1), sycl::read_constant); + __SYCL_SUPPRESS_DEPRECATED_POP sycl::accessor acc_7(buf_data, sycl::write_only, sycl::no_init); sycl::accessor acc_8(buf_data, sycl::range<1>(8), sycl::write_only, sycl::no_init); @@ -255,10 +265,13 @@ int main() { auto acc_2 = buf_data.get_access(sycl::range<1>(8), sycl::no_init); auto acc_3 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1), sycl::no_init); + // sycl::read_constant is deprecated. + __SYCL_SUPPRESS_DEPRECATED_PUSH auto acc_4 = buf_data.get_access(sycl::read_constant); auto acc_5 = buf_data.get_access(sycl::range<1>(8), sycl::read_constant); auto acc_6 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1), sycl::read_constant); + __SYCL_SUPPRESS_DEPRECATED_POP auto acc_7 = buf_data.get_access(sycl::write_only, sycl::no_init); auto acc_8 = buf_data.get_access(sycl::range<1>(8), sycl::write_only, sycl::no_init);