[SYCL] Implement multi_ptr default to be legacy to avoid code break with SYCL 1.2.1

sycl/include/sycl/access/access.hpp

@@ -56,7 +56,12 @@ enum class address_space : int {
   generic_space = 6, // TODO generic_space address space is not supported yet
 };
 
-enum class decorated : int { no = 0, yes = 1, legacy = 2 };
+enum class decorated : int {
+  no = 0,
+  yes = 1,
+  legacy __SYCL2020_DEPRECATED("sycl::access::decorated::legacy "
+                               "is deprecated since SYCL 2020") = 2
+};
 } // namespace access
 
 using access::target;

sycl/include/sycl/accessor.hpp

@@ -2292,12 +2292,22 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   template <access::target AccessTarget_ = AccessTarget,
             typename = std::enable_if_t<
                 (AccessTarget_ == access::target::host_buffer) ||
-                (AccessTarget_ == access::target::host_task) ||
-                (AccessTarget_ == access::target::device)>>
+                (AccessTarget_ == access::target::host_task)>>
   std::add_pointer_t<value_type> get_pointer() const noexcept {
     return getPointerAdjusted();
   }
 
+  template <
+      access::target AccessTarget_ = AccessTarget,
+      typename = std::enable_if_t<(AccessTarget_ == access::target::device)>>
+  __SYCL2020_DEPRECATED(
+      "accessor::get_pointer() is deprecated, please use get_multi_ptr()")
+  global_ptr<DataT> get_pointer() const noexcept {
+    return global_ptr<DataT>(
+        const_cast<typename detail::DecoratedType<DataT, AS>::type *>(
+            getPointerAdjusted()));
+  }
+
   template <access::target AccessTarget_ = AccessTarget,
             typename = std::enable_if_t<AccessTarget_ ==
                                         access::target::constant_buffer>>
@@ -3064,8 +3074,10 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
     return const_reverse_iterator(begin());
   }
 
-  std::add_pointer_t<value_type> get_pointer() const noexcept {
-    return std::add_pointer_t<value_type>(local_acc::getQualifiedPtr());
+  __SYCL2020_DEPRECATED(
+      "local_accessor::get_pointer() is deprecated, please use get_multi_ptr()")
+  local_ptr<DataT> get_pointer() const noexcept {
+    return local_ptr<DataT>(local_acc::getQualifiedPtr());
   }
 
   template <access::decorated IsDecorated>

sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp

@@ -82,7 +82,7 @@ using EnableIfAccessor = std::enable_if_t<
 template <typename T, int Dimensions>
 __ESIMD_API uint32_t localAccessorToOffset(local_accessor<T, Dimensions> acc) {
   return static_cast<uint32_t>(
-      reinterpret_cast<std::uintptr_t>(acc.get_pointer()));
+      reinterpret_cast<std::uintptr_t>(acc.get_pointer().get()));
 }
 
 } // namespace ext::intel::esimd::detail

sycl/include/sycl/ext/intel/esimd/detail/util.hpp

@@ -189,7 +189,8 @@ auto accessorToPointer(AccessorTy Acc, OffsetTy Offset = 0) {
   using QualTPtrType =
       std::conditional_t<std::is_const_v<typename AccessorTy::value_type>,
                          const T *, T *>;
-  auto BytePtr = reinterpret_cast<QualCharPtrType>(Acc.get_pointer()) + Offset;
+  auto BytePtr =
+      reinterpret_cast<QualCharPtrType>(Acc.get_pointer().get()) + Offset;
   return reinterpret_cast<QualTPtrType>(BytePtr);
 }
 #endif // __ESIMD_FORCE_STATELESS_MEM

sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp

@@ -729,7 +729,7 @@ lsc_gather(AccessorTy acc,
            __ESIMD_NS::simd_mask<N> pred = 1) {
 #ifdef __ESIMD_FORCE_STATELESS_MEM
   return lsc_gather<T, NElts, DS, L1H, L3H>(
-      reinterpret_cast<T *>(acc.get_pointer()), offsets, pred);
+      reinterpret_cast<T *>(acc.get_pointer().get()), offsets, pred);
 #else
   detail::check_lsc_vector_size<NElts>();
   detail::check_lsc_data_size<T, DS>();
@@ -803,7 +803,8 @@ lsc_gather(AccessorTy acc,
            __ESIMD_NS::simd<T, N * NElts> old_values) {
 #ifdef __ESIMD_FORCE_STATELESS_MEM
   return lsc_gather<T, NElts, DS, L1H, L3H>(
-      reinterpret_cast<T *>(acc.get_pointer()), offsets, pred, old_values);
+      reinterpret_cast<T *>(acc.get_pointer().get()), offsets, pred,
+      old_values);
 
 #else
   detail::check_lsc_vector_size<NElts>();

sycl/include/sycl/group.hpp

@@ -526,9 +526,11 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
   /// Permitted types for DestDataT are all scalar and vector types. SrcDataT
   /// must be either the same as DestDataT or const DestDataT.
   template <typename DestDataT, typename SrcDataT>
-  device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
-                                     decorated_global_ptr<SrcDataT> src,
-                                     size_t numElements) const {
+  typename std::enable_if_t<
+      std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
+  async_work_group_copy(decorated_local_ptr<DestDataT> dest,
+                        decorated_global_ptr<SrcDataT> src,
+                        size_t numElements) const {
     return async_work_group_copy(dest, src, numElements, 1);
   }
 
@@ -539,9 +541,11 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
   /// Permitted types for DestDataT are all scalar and vector types. SrcDataT
   /// must be either the same as DestDataT or const DestDataT.
   template <typename DestDataT, typename SrcDataT>
-  device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
-                                     decorated_local_ptr<SrcDataT> src,
-                                     size_t numElements) const {
+  typename std::enable_if_t<
+      std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
+  async_work_group_copy(decorated_global_ptr<DestDataT> dest,
+                        decorated_local_ptr<SrcDataT> src,
+                        size_t numElements) const {
     return async_work_group_copy(dest, src, numElements, 1);
   }
 

sycl/include/sycl/multi_ptr.hpp

@@ -662,6 +662,7 @@ class multi_ptr<void, Space, DecorateAddress> {
 template <typename ElementType, access::address_space Space>
 class multi_ptr<ElementType, Space, access::decorated::legacy> {
 public:
+  using value_type = ElementType;
   using element_type =
       std::conditional_t<std::is_same_v<ElementType, half>,
                          sycl::detail::half_impl::BIsRepresentationT,
@@ -773,9 +774,8 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
            Space == access::address_space::ext_intel_global_device_space)>>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
                      isPlaceholder, PropertyListT>
-                Accessor) {
-    m_Pointer = detail::cast_AS<pointer_t>(Accessor.get_pointer());
-  }
+                Accessor)
+      : multi_ptr(detail::cast_AS<pointer_t>(Accessor.get_pointer().get())) {}
 
   // Only if Space == local_space || generic_space
   template <
@@ -891,6 +891,10 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
 
   // Returns the underlying OpenCL C pointer
   pointer_t get() const { return m_Pointer; }
+  pointer_t get_decorated() const { return m_Pointer; }
+  std::add_pointer_t<element_type> get_raw() const {
+    return reinterpret_cast<std::add_pointer_t<element_type>>(get());
+  }
 
   // Implicit conversion to the underlying pointer type
   operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); }
@@ -1003,6 +1007,7 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
 template <access::address_space Space>
 class multi_ptr<void, Space, access::decorated::legacy> {
 public:
+  using value_type = void;
   using element_type = void;
   using difference_type = std::ptrdiff_t;
 
@@ -1114,6 +1119,10 @@ class multi_ptr<void, Space, access::decorated::legacy> {
   using ReturnPtr = detail::const_if_const_AS<Space, void> *;
   // Returns the underlying OpenCL C pointer
   pointer_t get() const { return m_Pointer; }
+  pointer_t get_decorated() const { return m_Pointer; }
+  std::add_pointer_t<element_type> get_raw() const {
+    return reinterpret_cast<std::add_pointer_t<element_type>>(get());
+  }
 
   // Implicit conversion to the underlying pointer type
   operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); };
@@ -1144,6 +1153,7 @@ class multi_ptr<void, Space, access::decorated::legacy> {
 template <access::address_space Space>
 class multi_ptr<const void, Space, access::decorated::legacy> {
 public:
+  using value_type = const void;
   using element_type = const void;
   using difference_type = std::ptrdiff_t;
 
@@ -1256,6 +1266,10 @@ class multi_ptr<const void, Space, access::decorated::legacy> {
 
   // Returns the underlying OpenCL C pointer
   pointer_t get() const { return m_Pointer; }
+  pointer_t get_decorated() const { return m_Pointer; }
+  std::add_pointer_t<element_type> get_raw() const {
+    return reinterpret_cast<std::add_pointer_t<element_type>>(get());
+  }
 
   // Implicit conversion to the underlying pointer type
   operator const void *() const {

sycl/include/sycl/nd_item.hpp

@@ -170,33 +170,39 @@ template <int Dimensions = 1> class nd_item {
   }
 
   template <typename DestDataT, typename SrcDataT>
-  device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
-                                     decorated_global_ptr<SrcDataT> src,
-                                     size_t numElements) const {
+  typename std::enable_if_t<
+      std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
+  async_work_group_copy(decorated_local_ptr<DestDataT> dest,
+                        decorated_global_ptr<SrcDataT> src,
+                        size_t numElements) const {
     return Group.async_work_group_copy(dest, src, numElements);
   }
 
   template <typename DestDataT, typename SrcDataT>
-  device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
-                                     decorated_local_ptr<SrcDataT> src,
-                                     size_t numElements) const {
+  typename std::enable_if_t<
+      std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
+  async_work_group_copy(decorated_global_ptr<DestDataT> dest,
+                        decorated_local_ptr<SrcDataT> src,
+                        size_t numElements) const {
     return Group.async_work_group_copy(dest, src, numElements);
   }
 
   template <typename DestDataT, typename SrcDataT>
-  device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
-                                     decorated_global_ptr<SrcDataT> src,
-                                     size_t numElements,
-                                     size_t srcStride) const {
+  typename std::enable_if_t<
+      std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
+  async_work_group_copy(decorated_local_ptr<DestDataT> dest,
+                        decorated_global_ptr<SrcDataT> src, size_t numElements,
+                        size_t srcStride) const {
 
     return Group.async_work_group_copy(dest, src, numElements, srcStride);
   }
 
   template <typename DestDataT, typename SrcDataT>
-  device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
-                                     decorated_local_ptr<SrcDataT> src,
-                                     size_t numElements,
-                                     size_t destStride) const {
+  typename std::enable_if_t<
+      std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
+  async_work_group_copy(decorated_global_ptr<DestDataT> dest,
+                        decorated_local_ptr<SrcDataT> src, size_t numElements,
+                        size_t destStride) const {
     return Group.async_work_group_copy(dest, src, numElements, destStride);
   }
 

sycl/test-e2e/ESIMD/accessor_local.cpp

@@ -53,7 +53,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
          uint32_t GID = Item.get_global_id(0);
          uint32_t LID = Item.get_local_id(0);
          uint32_t LocalAccOffset = static_cast<uint32_t>(
-             reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer()));
+             reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer().get()));
          if constexpr (TestSubscript) {
            for (int I = 0; I < VL; I++)
              LocalAcc[LID * VL + I] = GID * 100 + I;

sycl/test-e2e/InvokeSimd/Regression/slm_load_store.cpp

@@ -45,7 +45,7 @@ ESIMD_INLINE void slm_load_store_test(
 
   uint32_t LocalAccOffset =
       static_cast<uint32_t>(
-          reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer())) +
+          reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer().get())) +
       LAByteOffset;
   auto Local1 = esimd::slm_block_load<dtype, VL>(LocalAccOffset);
   auto Local2 = esimd::slm_block_load<dtype, VL>(LocalAccOffset +

sycl/test/basic_tests/accessor/accessor_get_pointer.cpp

@@ -27,8 +27,6 @@ void test_get_multi_ptr(handler &cgh, buffer<int, size> &buffer) {
   auto device_acc_ptr = device_acc.get_pointer();
   static_assert(std::is_same_v<decltype(acc_ptr), std::add_pointer_t<int>>);
   static_assert(std::is_same_v<decltype(target_local_ptr), local_ptr<int>>);
-  static_assert(
-      std::is_same_v<decltype(local_pointer), std::add_pointer_t<int>>);
-  static_assert(
-      std::is_same_v<decltype(device_acc_ptr), std::add_pointer_t<int>>);
+  static_assert(std::is_same_v<decltype(local_pointer), local_ptr<int>>);
+  static_assert(std::is_same_v<decltype(device_acc_ptr), global_ptr<int>>);
 }

sycl/test/esimd/esimd_verify.cpp

@@ -9,7 +9,7 @@
 using namespace sycl;
 using namespace sycl::ext::intel::esimd;
 
-// CHECK-NEGATIVE-DAG: error: function 'int* sycl::_V1::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
+// CHECK-NEGATIVE-DAG: error: function 'sycl::_V1::multi_ptr<{{.+}}> sycl::_V1::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
 // CHECK-NEGATIVE-DAG: error: function '{{.+}} sycl::_V1::accessor<{{.+}}>::operator[]<{{.+}}>({{.+}}) const' is not supported in ESIMD context
 // CHECK-NEGATIVE-DAG: error: function '{{.+}}combine(int const&)' is not supported in ESIMD context
 

sycl/test/extensions/usm_pointers_aliases.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify %s -o %t.out
+// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify %s -o %t.out -Wno-deprecated-declarations
 // expected-no-diagnostics
 
 #include <sycl/sycl.hpp>

sycl/test/multi_ptr/ctad.cpp

@@ -49,8 +49,11 @@ int main() {
                                          sycl::access::decorated::no>;
   using constMPtr = sycl::multi_ptr<int, address_space::constant_space,
                                     sycl::access::decorated::legacy>;
+  using constDefaultMPtr = sycl::multi_ptr<int, address_space::constant_space>;
   using localMPtr = sycl::multi_ptr<int, address_space::local_space,
                                     sycl::access::decorated::no>;
+  using legacyMPtr = sycl::multi_ptr<int, address_space::global_space,
+                                     sycl::access::decorated::legacy>;
   static_assert(std::is_same<rwDeviceCTAD, deviceMPtr>::value);
   static_assert(std::is_same<rwDeviceCTAD, globlMPtr>::value);
   static_assert(std::is_same<rwGloblCTAD, globlMPtr>::value);
@@ -63,6 +66,15 @@ int main() {
   static_assert(std::is_same<constCTAD, constMPtr>::value);
   static_assert(std::is_same<localCTAD, localMPtr>::value);
   static_assert(std::is_same<localCTADDep, localMPtr>::value);
+  static_assert(std::is_same<constMPtr, constDefaultMPtr>::value);
+
+  legacyMPtr LegacytMultiPtr;
+  static_assert(
+      std::is_same_v<
+          decltype(LegacytMultiPtr.get_decorated()),
+          typename sycl::multi_ptr<int, address_space::global_space,
+                                   sycl::access::decorated::yes>::pointer>);
+  static_assert(std::is_same_v<decltype(LegacytMultiPtr.get_raw()), int *>);
 
   globlMPtr non_const_multi_ptr;
   auto constTypeMultiPtr = constTypeMPtr(non_const_multi_ptr);

sycl/test/warnings/sycl_2020_deprecations.cpp

@@ -303,21 +303,31 @@ int main() {
         [=](sycl::nd_item<1> Idx) {
           int PrivateVal = 0;
 
+          // expected-warning@+6{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
+          // expected-warning@+8{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
+          // expected-warning@+8{{'get_pointer' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
+          // expected-warning@+7{{'get_pointer<sycl::access::target::global_buffer, void>' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
           // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::global_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
           sycl::multi_ptr<int, sycl::access::address_space::global_space,
                           sycl::access::decorated::legacy>
               LegacyGlobalMptr =
                   sycl::make_ptr<int, sycl::access::address_space::global_space,
                                  sycl::access::decorated::legacy>(
                       GlobalAcc.get_pointer());
+          // expected-warning@+5{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
+          // expected-warning@+7{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
+          // expected-warning@+7{{'get_pointer' is deprecated: local_accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
           // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
           sycl::multi_ptr<int, sycl::access::address_space::local_space,
                           sycl::access::decorated::legacy>
               LegacyLocalMptr =
                   sycl::make_ptr<int, sycl::access::address_space::local_space,
                                  sycl::access::decorated::legacy>(
                       LocalAcc.get_pointer());
-          // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
+
+          // expected-warning@+4{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
+          // expected-warning@+5{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
+          // expected-warning@+6{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
           sycl::multi_ptr<int, sycl::access::address_space::private_space,
                           sycl::access::decorated::legacy>
               LegacyPrivateMptr =