[SYCL] Enabling depreacate warnings for sycl

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)

sycl/include/sycl/accessor.hpp

@@ -362,9 +362,9 @@ class accessor_common {
       return MAccessor[MIDs];
     }
 
-    template <int CurDims = SubDims>
-    typename std::enable_if_t<CurDims == 1 && IsAccessAtomic, atomic<DataT, AS>>
-    operator[](size_t Index) const {
+    template <int CurDims = SubDims,
+              typename = std::enable_if_t<CurDims == 1 && IsAccessAtomic>>
+    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 <int Dims = Dimensions>
-  operator typename std::enable_if_t<Dims == 0 &&
-                                         AccessMode == access::mode::atomic,
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<Dims == 0 &&
+                                        AccessMode == access::mode::atomic>>
+  operator
 #ifdef __ENABLE_USM_ADDR_SPACE__
-                                     atomic<DataT>
+      atomic<DataT>
 #else
-                                     atomic<DataT, AS>
+      atomic<DataT, AS>
 #endif
-                                     >() const {
+      () const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
-    return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
-        getQualifiedPtr() + LinearIndex));
+    return multi_ptr<DataT, AS, access::decorated::yes>(getQualifiedPtr() +
+                                                        LinearIndex);
   }
 
-  template <int Dims = Dimensions>
-  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
-                            atomic<DataT, AS>>
-  operator[](id<Dimensions> Index) const {
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<(Dims > 0) &&
+                                        AccessMode == access::mode::atomic>>
+  auto operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
         getQualifiedPtr() + LinearIndex));
   }
 
-  template <int Dims = Dimensions>
-  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
-                            atomic<DataT, AS>>
-  operator[](size_t Index) const {
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<Dims == 1 &&
+                                        AccessMode == access::mode::atomic>>
+  auto operator[](size_t Index) const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
     return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
         getQualifiedPtr() + LinearIndex));
   }
   template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
-  auto operator[](size_t Index) const {
+  AccessorSubscript<Dims - 1> operator[](size_t Index) const {
     return AccessorSubscript<Dims - 1>(*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<AdjustedDim, 1>(getMemoryRange()),
         getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
 
   iterator end() const noexcept {
     return iterator::getEnd(
-        get_pointer(),
+        getPointerAdjusted(),
         detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
         getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
 
   const_iterator cbegin() const noexcept {
     return const_iterator::getBegin(
-        get_pointer(),
+        getPointerAdjusted(),
         detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
         getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
 
   const_iterator cend() const noexcept {
     return const_iterator::getEnd(
-        get_pointer(),
+        getPointerAdjusted(),
         detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
         getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
@@ -2344,27 +2345,27 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
     return getQualifiedPtr()[Index];
   }
 
-  template <int Dims = Dimensions>
-  operator typename std::enable_if_t<
-      Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
-      const {
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<Dims == 0 &&
+                                        AccessMode == access::mode::atomic>>
+  operator atomic<DataT, AS>() const {
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS, access::decorated::yes>(getQualifiedPtr()));
   }
 
-  template <int Dims = Dimensions>
-  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
-                            atomic<DataT, AS>>
-  operator[](id<Dimensions> Index) const {
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<(Dims > 0) &&
+                                        AccessMode == access::mode::atomic>>
+  auto operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
         getQualifiedPtr() + LinearIndex));
   }
 
-  template <int Dims = Dimensions>
-  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
-                            atomic<DataT, AS>>
-  operator[](size_t Index) const {
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<Dims == 1 &&
+                                        AccessMode == access::mode::atomic>>
+  auto operator[](size_t Index) const {
     return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
         getQualifiedPtr() + Index));
   }

sycl/include/sycl/atomic.hpp

@@ -169,8 +169,7 @@ inline namespace _V1 {
 
 template <typename T, access::address_space addressSpace =
                           access::address_space::global_space>
-class __SYCL2020_DEPRECATED(
-    "sycl::atomic is deprecated since SYCL 2020") atomic {
+class atomic {
   friend class atomic<T, access::address_space::global_space>;
   static_assert(detail::IsValidAtomicType<T>::value,
                 "Invalid SYCL atomic type. Valid types are: int, "
@@ -194,9 +193,11 @@ class __SYCL2020_DEPRECATED(
 public:
   template <typename pointerT, access::decorated IsDecorated>
 #ifdef __SYCL_DEVICE_ONLY__
+  __SYCL2020_DEPRECATED("use sycl::atomic_ref instead")
   atomic(multi_ptr<pointerT, addressSpace, IsDecorated> ptr)
       : Ptr(GetDecoratedPtr(ptr))
 #else
+  __SYCL2020_DEPRECATED("use sycl::atomic_ref instead")
   atomic(multi_ptr<pointerT, addressSpace, IsDecorated> ptr)
       : Ptr(reinterpret_cast<std::atomic<T> *>(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<T, access::address_space::ext_intel_global_device_space>
              &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<T, access::address_space::ext_intel_global_device_space> &&RHS) {
     Ptr = RHS.Ptr;
@@ -400,7 +403,6 @@ T atomic_fetch_max(atomic<T, addressSpace> Object, T Operand,
                    memory_order MemoryOrder = memory_order::relaxed) {
   return Object.fetch_max(Operand, MemoryOrder);
 }
-
 } // namespace _V1
 } // namespace sycl
 

sycl/include/sycl/detail/is_device_copyable.hpp

@@ -100,14 +100,24 @@ struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
     IsDeprecatedDeviceCopyable<T[N]> : IsDeprecatedDeviceCopyable<T> {};
 
 #ifdef __SYCL_DEVICE_ONLY__
+// Helper trait to avoid instantiated the deprecated specializations of
+// IsDeprecatedDeviceCopyable when is_device_copyable_v<T> is true.
+template <typename T, typename = void>
+struct IsDeprecatedDeviceCopyableIfNot2020 : std::false_type {};
+
+template <typename T>
+struct IsDeprecatedDeviceCopyableIfNot2020<
+    T, std::enable_if_t<!is_device_copyable_v<T>>>
+    : detail::IsDeprecatedDeviceCopyable<T> {};
+
 // Checks that the fields of the type T with indices 0 to (NumFieldsToCheck -
 // 1) are device copyable.
 template <typename T, unsigned NumFieldsToCheck>
 struct CheckFieldsAreDeviceCopyable
     : CheckFieldsAreDeviceCopyable<T, NumFieldsToCheck - 1> {
   using FieldT = decltype(__builtin_field_type(T, NumFieldsToCheck - 1));
   static_assert(is_device_copyable_v<FieldT> ||
-                    detail::IsDeprecatedDeviceCopyable<FieldT>::value,
+                    detail::IsDeprecatedDeviceCopyableIfNot2020<FieldT>::value,
                 "The specified type is not device copyable");
 };
 
@@ -120,7 +130,7 @@ struct CheckBasesAreDeviceCopyable
     : CheckBasesAreDeviceCopyable<T, NumBasesToCheck - 1> {
   using BaseT = decltype(__builtin_base_type(T, NumBasesToCheck - 1));
   static_assert(is_device_copyable_v<BaseT> ||
-                    detail::IsDeprecatedDeviceCopyable<BaseT>::value,
+                    detail::IsDeprecatedDeviceCopyableIfNot2020<BaseT>::value,
                 "The specified type is not device copyable");
 };
 

sycl/include/sycl/ext/oneapi/experimental/root_group.hpp

@@ -101,7 +101,8 @@ void group_barrier(ext::oneapi::experimental::root_group<dimensions> 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<dimensions>();
+  const auto ChildGroup =
+      ext::oneapi::this_work_item::get_work_group<dimensions>();
   if (ChildGroup.get_group_linear_range() == 1) {
     group_barrier(ChildGroup);
   } else {

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-=//

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<const InputT>(in),
-                                  global_ptr<const InputT>(in) + N, binary_op);
-            out[3] =
-                joint_reduce(g, global_ptr<const InputT>(in),
-                             global_ptr<const InputT>(in) + N, init, binary_op);
-            out[4] = joint_reduce(sg, global_ptr<const InputT>(in),
-                                  global_ptr<const InputT>(in) + N, binary_op);
-            out[5] =
-                joint_reduce(sg, global_ptr<const InputT>(in),
-                             global_ptr<const InputT>(in) + N, init, binary_op);
+            out[2] = joint_reduce(g, decorated_global_ptr<const InputT>(in),
+                                  decorated_global_ptr<const InputT>(in) + N,
+                                  binary_op);
+            out[3] = joint_reduce(g, decorated_global_ptr<const InputT>(in),
+                                  decorated_global_ptr<const InputT>(in) + N,
+                                  init, binary_op);
+            out[4] = joint_reduce(sg, decorated_global_ptr<const InputT>(in),
+                                  decorated_global_ptr<const InputT>(in) + N,
+                                  binary_op);
+            out[5] = joint_reduce(sg, decorated_global_ptr<const InputT>(in),
+                                  decorated_global_ptr<const InputT>(in) + N,
+                                  init, binary_op);
           });
     });
   }

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

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

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<const volatile sycl::int2>(&global[i]);
             auto x_cv2 = sg.load(
-                sycl::global_ptr<const volatile sycl::int2>(&global[i]));
-
+                global.get_multi_ptr<sycl::access::decorated::yes>() + i);
             // Local address space
             auto y = sg.load(&local[i]);
             auto y_cv1 = sg.load<const volatile sycl::int2>(&local[i]);
-            auto y_cv2 =
-                sg.load(sycl::local_ptr<const volatile sycl::int2>(&local[i]));
+            auto y_cv2 = sg.load(
+                local.get_multi_ptr<sycl::access::decorated::yes>() + i);
 
             // Store result only if same for non-cv and cv
             if (utils<int, 2>::cmp_vec(x, x_cv1) &&

sycl/test/warnings/sycl_2020_deprecations.cpp

@@ -150,7 +150,7 @@ int main() {
   sycl::multi_ptr<int, sycl::access::address_space::global_space,
                   sycl::access::decorated::yes>
       a(nullptr);
-  // expected-warning@+1 {{'atomic<int, sycl::access::address_space::global_space>' is deprecated: sycl::atomic is deprecated since SYCL 2020}}
+  // expected-warning@+1 {{'atomic' is deprecated: use sycl::atomic_ref instead}}
   sycl::atomic<int> 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<int, 1, sycl::access::mode::read_write, sycl::target::local>
         LocalAcc(sycl::range<1>(1), CGH);
+    // expected-warning@sycl/pointers.hpp:33 {{'multi_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy>' 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<int, sycl::access::address_space::local_space,
                     sycl::access::decorated::no>
@@ -261,6 +262,7 @@ int main() {
     // expected-warning@+1{{'local' is deprecated: use `local_accessor` instead}}
     sycl::accessor<const int, 1, sycl::access::mode::read, sycl::target::local>
         LocalConstAcc(sycl::range<1>(1), CGH);
+    // expected-warning@sycl/pointers.hpp:33 {{'multi_ptr<const int, sycl::access::address_space::local_space, sycl::access::decorated::legacy>' 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<const int, sycl::access::address_space::local_space,
                     sycl::access::decorated::no>