diff --git a/clang/lib/Sema/SemaAvailability.cpp b/clang/lib/Sema/SemaAvailability.cpp index f3c462f17cb75..e04cbeec16555 100644 --- a/clang/lib/Sema/SemaAvailability.cpp +++ b/clang/lib/Sema/SemaAvailability.cpp @@ -640,8 +640,7 @@ static void DoEmitAvailabilityWarning(Sema &S, AvailabilityResult K, S.getTopMostPointOfInstantiation(ReferringDecl); bool ShouldAllowWarningInSystemHeader = InstantiationLoc != Loc && - !S.getSourceManager().isInSystemHeader(InstantiationLoc) && - !S.getLangOpts().SYCLIsDevice && !S.getLangOpts().SYCLIsHost; + !S.getSourceManager().isInSystemHeader(InstantiationLoc); struct AllowWarningInSystemHeaders { AllowWarningInSystemHeaders(DiagnosticsEngine &E, bool AllowWarningInSystemHeaders) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 5ff42e386699a..de48eaa20cbf8 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -362,9 +362,9 @@ class accessor_common { return MAccessor[MIDs]; } - template - typename std::enable_if_t> - operator[](size_t Index) const { + template > + auto operator[](size_t Index) const { MIDs[Dims - CurDims] = Index; return MAccessor[MIDs]; } @@ -1726,39 +1726,40 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : return getQualifiedPtr()[LinearIndex]; } - template - operator typename std::enable_if_t> + operator #ifdef __ENABLE_USM_ADDR_SPACE__ - atomic + atomic #else - atomic + atomic #endif - >() const { + () const { const size_t LinearIndex = getLinearIndex(id()); - return atomic(multi_ptr( - getQualifiedPtr() + LinearIndex)); + return multi_ptr(getQualifiedPtr() + + LinearIndex); } - template - typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic, - atomic> - operator[](id Index) const { + template 0) && + AccessMode == access::mode::atomic>> + auto operator[](id Index) const { const size_t LinearIndex = getLinearIndex(Index); return atomic(multi_ptr( getQualifiedPtr() + LinearIndex)); } - template - typename std::enable_if_t> - operator[](size_t Index) const { + template > + auto operator[](size_t Index) const { const size_t LinearIndex = getLinearIndex(id(Index)); return atomic(multi_ptr( getQualifiedPtr() + LinearIndex)); } template 1)>> - auto operator[](size_t Index) const { + AccessorSubscript operator[](size_t Index) const { return AccessorSubscript(*this, Index); } @@ -1851,28 +1852,28 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : iterator begin() const noexcept { return iterator::getBegin( - get_pointer(), + getPointerAdjusted(), detail::convertToArrayOfN(getMemoryRange()), getRange(), getOffset()); } iterator end() const noexcept { return iterator::getEnd( - get_pointer(), + getPointerAdjusted(), detail::convertToArrayOfN(getMemoryRange()), getRange(), getOffset()); } const_iterator cbegin() const noexcept { return const_iterator::getBegin( - get_pointer(), + getPointerAdjusted(), detail::convertToArrayOfN(getMemoryRange()), getRange(), getOffset()); } const_iterator cend() const noexcept { return const_iterator::getEnd( - get_pointer(), + getPointerAdjusted(), detail::convertToArrayOfN(getMemoryRange()), getRange(), getOffset()); } @@ -2344,27 +2345,27 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : return getQualifiedPtr()[Index]; } - template - operator typename std::enable_if_t< - Dims == 0 && AccessMode == access::mode::atomic, atomic>() - const { + template > + operator atomic() const { return atomic( multi_ptr(getQualifiedPtr())); } - template - typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic, - atomic> - operator[](id Index) const { + template 0) && + AccessMode == access::mode::atomic>> + auto operator[](id Index) const { const size_t LinearIndex = getLinearIndex(Index); return atomic(multi_ptr( getQualifiedPtr() + LinearIndex)); } - template - typename std::enable_if_t> - operator[](size_t Index) const { + template > + auto operator[](size_t Index) const { return atomic(multi_ptr( getQualifiedPtr() + Index)); } diff --git a/sycl/include/sycl/atomic.hpp b/sycl/include/sycl/atomic.hpp index 8cae0c047c0a1..361462e3cda89 100644 --- a/sycl/include/sycl/atomic.hpp +++ b/sycl/include/sycl/atomic.hpp @@ -169,8 +169,7 @@ inline namespace _V1 { template -class __SYCL2020_DEPRECATED( - "sycl::atomic is deprecated since SYCL 2020") atomic { +class atomic { friend class atomic; static_assert(detail::IsValidAtomicType::value, "Invalid SYCL atomic type. Valid types are: int, " @@ -194,9 +193,11 @@ class __SYCL2020_DEPRECATED( public: template #ifdef __SYCL_DEVICE_ONLY__ + __SYCL2020_DEPRECATED("use sycl::atomic_ref instead") atomic(multi_ptr ptr) : Ptr(GetDecoratedPtr(ptr)) #else + __SYCL2020_DEPRECATED("use sycl::atomic_ref instead") atomic(multi_ptr ptr) : Ptr(reinterpret_cast *>(ptr.get())) #endif @@ -211,6 +212,7 @@ class __SYCL2020_DEPRECATED( typename = typename std::enable_if_t< _Space == addressSpace && addressSpace == access::address_space::global_space>> + __SYCL2020_DEPRECATED("use sycl::atomic_ref instead") atomic(const atomic &RHS) { Ptr = RHS.Ptr; @@ -220,6 +222,7 @@ class __SYCL2020_DEPRECATED( typename = typename std::enable_if_t< _Space == addressSpace && addressSpace == access::address_space::global_space>> + __SYCL2020_DEPRECATED("use sycl::atomic_ref instead") atomic( atomic &&RHS) { Ptr = RHS.Ptr; @@ -400,7 +403,6 @@ T atomic_fetch_max(atomic Object, T Operand, memory_order MemoryOrder = memory_order::relaxed) { return Object.fetch_max(Operand, MemoryOrder); } - } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/detail/is_device_copyable.hpp b/sycl/include/sycl/detail/is_device_copyable.hpp index 388029e6a16a3..3b553b09a0e46 100644 --- a/sycl/include/sycl/detail/is_device_copyable.hpp +++ b/sycl/include/sycl/detail/is_device_copyable.hpp @@ -100,6 +100,16 @@ struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020") IsDeprecatedDeviceCopyable : IsDeprecatedDeviceCopyable {}; #ifdef __SYCL_DEVICE_ONLY__ +// Helper trait to avoid instantiated the deprecated specializations of +// IsDeprecatedDeviceCopyable when is_device_copyable_v is true. +template +struct IsDeprecatedDeviceCopyableIfNot2020 : std::false_type {}; + +template +struct IsDeprecatedDeviceCopyableIfNot2020< + T, std::enable_if_t>> + : detail::IsDeprecatedDeviceCopyable {}; + // Checks that the fields of the type T with indices 0 to (NumFieldsToCheck - // 1) are device copyable. template @@ -107,7 +117,7 @@ struct CheckFieldsAreDeviceCopyable : CheckFieldsAreDeviceCopyable { using FieldT = decltype(__builtin_field_type(T, NumFieldsToCheck - 1)); static_assert(is_device_copyable_v || - detail::IsDeprecatedDeviceCopyable::value, + detail::IsDeprecatedDeviceCopyableIfNot2020::value, "The specified type is not device copyable"); }; @@ -120,7 +130,7 @@ struct CheckBasesAreDeviceCopyable : CheckBasesAreDeviceCopyable { using BaseT = decltype(__builtin_base_type(T, NumBasesToCheck - 1)); static_assert(is_device_copyable_v || - detail::IsDeprecatedDeviceCopyable::value, + detail::IsDeprecatedDeviceCopyableIfNot2020::value, "The specified type is not device copyable"); }; diff --git a/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp index b8c90683bbaaf..40ec43c4bfa20 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp @@ -101,7 +101,8 @@ void group_barrier(ext::oneapi::experimental::root_group G, // workaround that's not intended to reduce the bar for SPIR-V modules // acceptance, but rather make a pessimistic case work until we have full // support for the device barrier built-in from backends. - const auto ChildGroup = ext::oneapi::experimental::this_group(); + const auto ChildGroup = + ext::oneapi::this_work_item::get_work_group(); if (ChildGroup.get_group_linear_range() == 1) { group_barrier(ChildGroup); } else { diff --git a/sycl/test-e2e/Basic/queue/queue_parallel_for_generic.cpp b/sycl/test-e2e/Basic/queue/queue_parallel_for_generic.cpp index eb00aa1038c21..bd4bba0355905 100644 --- a/sycl/test-e2e/Basic/queue/queue_parallel_for_generic.cpp +++ b/sycl/test-e2e/Basic/queue/queue_parallel_for_generic.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out +// RUN: %{build} -Wno-error=deprecated-declarations -o %t.out // RUN: %{run} %t.out //=-queue_parallel_for_generic.cpp - SYCL queue parallel_for generic lambda-=// diff --git a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp index dce0ee222d864..bd994209e5f99 100644 --- a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp +++ b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp @@ -44,16 +44,18 @@ void test(const InputContainer &input, BinaryOperation binary_op, int lid = it.get_local_id(0); out[0] = reduce_over_group(g, in[lid], binary_op); out[1] = reduce_over_group(g, in[lid], init, binary_op); - out[2] = joint_reduce(g, global_ptr(in), - global_ptr(in) + N, binary_op); - out[3] = - joint_reduce(g, global_ptr(in), - global_ptr(in) + N, init, binary_op); - out[4] = joint_reduce(sg, global_ptr(in), - global_ptr(in) + N, binary_op); - out[5] = - joint_reduce(sg, global_ptr(in), - global_ptr(in) + N, init, binary_op); + out[2] = joint_reduce(g, decorated_global_ptr(in), + decorated_global_ptr(in) + N, + binary_op); + out[3] = joint_reduce(g, decorated_global_ptr(in), + decorated_global_ptr(in) + N, + init, binary_op); + out[4] = joint_reduce(sg, decorated_global_ptr(in), + decorated_global_ptr(in) + N, + binary_op); + out[5] = joint_reduce(sg, decorated_global_ptr(in), + decorated_global_ptr(in) + N, + init, binary_op); }); }); } diff --git a/sycl/test-e2e/InorderQueue/in_order_event_release.cpp b/sycl/test-e2e/InorderQueue/in_order_event_release.cpp index 47a6c42a8ea43..1d3c9207ca04e 100644 --- a/sycl/test-e2e/InorderQueue/in_order_event_release.cpp +++ b/sycl/test-e2e/InorderQueue/in_order_event_release.cpp @@ -1,5 +1,5 @@ // -// RUN: %{build} -o %t.out +// RUN: %{build} -Wno-error=deprecated-declarations -o %t.out // RUN: %{run} %t.out // SYCL ordered queue event release shortcut test diff --git a/sycl/test-e2e/InorderQueue/in_order_kernels.cpp b/sycl/test-e2e/InorderQueue/in_order_kernels.cpp index daa2f553a1ef9..3a525b78009f8 100644 --- a/sycl/test-e2e/InorderQueue/in_order_kernels.cpp +++ b/sycl/test-e2e/InorderQueue/in_order_kernels.cpp @@ -1,5 +1,5 @@ // -// RUN: %{build} -o %t.out +// RUN: %{build} -Wno-error=deprecated-declarations -o %t.out // RUN: %{run} %t.out // SYCL ordered queue kernel shortcut test diff --git a/sycl/test-e2e/SubGroup/sub_group_as_vec.cpp b/sycl/test-e2e/SubGroup/sub_group_as_vec.cpp index 210c147d6a776..b37f672408b76 100644 --- a/sycl/test-e2e/SubGroup/sub_group_as_vec.cpp +++ b/sycl/test-e2e/SubGroup/sub_group_as_vec.cpp @@ -54,13 +54,12 @@ int main(int argc, char *argv[]) { auto x = sg.load(&global[i]); auto x_cv1 = sg.load(&global[i]); auto x_cv2 = sg.load( - sycl::global_ptr(&global[i])); - + global.get_multi_ptr() + i); // Local address space auto y = sg.load(&local[i]); auto y_cv1 = sg.load(&local[i]); - auto y_cv2 = - sg.load(sycl::local_ptr(&local[i])); + auto y_cv2 = sg.load( + local.get_multi_ptr() + i); // Store result only if same for non-cv and cv if (utils::cmp_vec(x, x_cv1) && diff --git a/sycl/test/warnings/sycl_2020_deprecations.cpp b/sycl/test/warnings/sycl_2020_deprecations.cpp index 4614707606830..49be1ef105238 100644 --- a/sycl/test/warnings/sycl_2020_deprecations.cpp +++ b/sycl/test/warnings/sycl_2020_deprecations.cpp @@ -150,7 +150,7 @@ int main() { sycl::multi_ptr a(nullptr); - // expected-warning@+1 {{'atomic' is deprecated: sycl::atomic is deprecated since SYCL 2020}} + // expected-warning@+1 {{'atomic' is deprecated: use sycl::atomic_ref instead}} sycl::atomic b(a); sycl::group<1> group = sycl::detail::Builder::createGroup<1>({8}, {4}, {1}); @@ -250,6 +250,7 @@ int main() { // expected-warning@+1{{'local' is deprecated: use `local_accessor` instead}} sycl::accessor LocalAcc(sycl::range<1>(1), CGH); + // expected-warning@sycl/pointers.hpp:33 {{'multi_ptr' is deprecated: decorated::legacy multi_ptr specialization is deprecated since SYCL 2020}} // expected-warning@+3{{'multi_ptr' is deprecated: multi_ptr construction using target::local specialized accessor is deprecated since SYCL 2020}} sycl::multi_ptr @@ -261,6 +262,7 @@ int main() { // expected-warning@+1{{'local' is deprecated: use `local_accessor` instead}} sycl::accessor LocalConstAcc(sycl::range<1>(1), CGH); + // expected-warning@sycl/pointers.hpp:33 {{'multi_ptr' is deprecated: decorated::legacy multi_ptr specialization is deprecated since SYCL 2020}} // expected-warning@+3{{'multi_ptr' is deprecated: multi_ptr construction using target::local specialized accessor is deprecated since SYCL 2020}} sycl::multi_ptr