[SYCL] Update get_pointer to return T* for target::device specialized accessor. (#8874)

* Update get_pointer to return T* for target::device specialized
accessor according to the
[Specification](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_buffer_accessor_for_commands)
* Declare src of `async_work_group_copy` in `group.hpp` and
`nd_item.hpp` as `global_ptr<const>`
* Implement implicit conversion from `multi_ptr<T>` to `multi_ptr<const
T>`
* Modifies `is_native_op` to cover both `const` and `non-const` types.

---------

Co-authored-by: Steffen Larsen <[email protected]>

Author: mmoadeli

sycl/include/sycl/accessor.hpp

@@ -2124,23 +2124,12 @@ 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)>>
-#if SYCL_LANGUAGE_VERSION >= 202001
-  std::add_pointer_t<value_type> get_pointer() const noexcept
-#else
-  DataT *get_pointer() const
-#endif
-  {
+                (AccessTarget_ == access::target::host_task) ||
+                (AccessTarget_ == access::target::device)>>
+  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>>
-  global_ptr<DataT> get_pointer() const {
-    return global_ptr<DataT>(getPointerAdjusted());
-  }
-
   template <access::target AccessTarget_ = AccessTarget,
             typename = std::enable_if_t<AccessTarget_ ==
                                         access::target::constant_buffer>>

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

@@ -182,9 +182,15 @@ template <unsigned N> class ForHelper {
 /// Returns the address referenced by the accessor \p Acc and
 /// the byte offset \p Offset.
 template <typename T, typename AccessorTy, typename OffsetTy = uint32_t>
-T *accessorToPointer(AccessorTy Acc, OffsetTy Offset = 0) {
-  auto BytePtr = reinterpret_cast<char *>(Acc.get_pointer().get()) + Offset;
-  return reinterpret_cast<T *>(BytePtr);
+auto accessorToPointer(AccessorTy Acc, OffsetTy Offset = 0) {
+  using QualCharPtrType =
+      std::conditional_t<std::is_const_v<typename AccessorTy::value_type>,
+                         const char *, char *>;
+  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;
+  return reinterpret_cast<QualTPtrType>(BytePtr);
 }
 #endif // __ESIMD_FORCE_STATELESS_MEM
 

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

@@ -776,8 +776,7 @@ __ESIMD_API std::enable_if_t<!std::is_pointer_v<AccessorTy>,
 lsc_gather(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
            __ESIMD_NS::simd_mask<N> pred = 1) {
 #ifdef __ESIMD_FORCE_STATELESS_MEM
-  return lsc_gather<T, NElts, DS, L1H, L3H>(acc.get_pointer().get(), offsets,
-                                            pred);
+  return lsc_gather<T, NElts, DS, L1H, L3H>(acc.get_pointer(), offsets, pred);
 #else
   detail::check_lsc_vector_size<NElts>();
   detail::check_lsc_data_size<T, DS>();
@@ -829,8 +828,8 @@ lsc_gather(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
            __ESIMD_NS::simd_mask<N> pred,
            __ESIMD_NS::simd<T, N * NElts> old_values) {
 #ifdef __ESIMD_FORCE_STATELESS_MEM
-  return lsc_gather<T, NElts, DS, L1H, L3H>(acc.get_pointer().get(), offsets,
-                                            pred, old_values);
+  return lsc_gather<T, NElts, DS, L1H, L3H>(acc.get_pointer(), offsets, pred,
+                                            old_values);
 #else
   detail::check_lsc_vector_size<NElts>();
   detail::check_lsc_data_size<T, DS>();

sycl/include/sycl/group_algorithm.hpp

@@ -96,7 +96,10 @@ using native_op_list =
 
 template <typename T, typename BinaryOperation> struct is_native_op {
   static constexpr bool value =
-      is_contained<BinaryOperation, native_op_list<T>>::value ||
+      is_contained<BinaryOperation,
+                   native_op_list<std::remove_const_t<T>>>::value ||
+      is_contained<BinaryOperation,
+                   native_op_list<std::add_const_t<T>>>::value ||
       is_contained<BinaryOperation, native_op_list<void>>::value;
 };
 
@@ -123,9 +126,9 @@ struct is_complex
 
 // ---- is_arithmetic_or_complex
 template <typename T>
-using is_arithmetic_or_complex =
-    std::integral_constant<bool, sycl::detail::is_complex<T>::value ||
-                                     sycl::detail::is_arithmetic<T>::value>;
+using is_arithmetic_or_complex = std::integral_constant<
+    bool, sycl::detail::is_complex<typename std::remove_cv_t<T>>::value ||
+              sycl::detail::is_arithmetic<T>::value>;
 
 template <typename T>
 struct is_vector_arithmetic_or_complex

sycl/include/sycl/multi_ptr.hpp

@@ -113,6 +113,16 @@ class multi_ptr {
       : m_Pointer(ptr) {}
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
 
+  // Implicit conversion from multi_ptr<T> to multi_ptr<const T>
+  template <typename NonConstElementType = std::remove_const_t<ElementType>,
+            typename = typename std::enable_if_t<
+                std::is_const_v<ElementType> &&
+                std::is_same_v<NonConstElementType,
+                               std::remove_const_t<ElementType>>>>
+  explicit multi_ptr(
+      multi_ptr<NonConstElementType, Space, DecorateAddress> MPtr)
+      : m_Pointer(MPtr.get_decorated()) {}
+
   // Only if Space is in
   // {global_space, ext_intel_global_device_space, generic_space}
   template <
@@ -126,8 +136,7 @@ class multi_ptr {
   multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(
-            detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}
+      : multi_ptr(Accessor.template get_multi_ptr<DecorateAddress>()) {}
 
   // Only if Space == local_space || generic_space
   template <int Dimensions, access::mode Mode,
@@ -149,7 +158,7 @@ class multi_ptr {
                 (Space == access::address_space::generic_space ||
                  Space == access::address_space::local_space)>>
   multi_ptr(local_accessor<ElementType, Dimensions> Accessor)
-      : m_Pointer(detail::cast_AS<decorated_type *>(Accessor.get_pointer())) {}
+      : multi_ptr(Accessor.template get_multi_ptr<DecorateAddress>()) {}
 
   // The following constructors are necessary to create multi_ptr<const
   // ElementType, Space, DecorateAddress> from accessor<ElementType, ...>.
@@ -177,8 +186,8 @@ class multi_ptr {
   multi_ptr(accessor<typename std::remove_const_t<RelayElementType>, Dimensions,
                      Mode, access::target::device, isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(
-            detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}
+      : m_Pointer(Accessor.template get_multi_ptr<DecorateAddress>()
+                      .get_decorated()) {}
 
   // Only if Space == local_space || generic_space and element type is const
   template <int Dimensions, access::mode Mode,
@@ -208,7 +217,7 @@ class multi_ptr {
   multi_ptr(
       local_accessor<typename std::remove_const_t<RelayElementType>, Dimensions>
           Accessor)
-      : m_Pointer(detail::cast_AS<decorated_type *>(Accessor.get_pointer())) {}
+      : multi_ptr(Accessor.template get_multi_ptr<DecorateAddress>()) {}
 
   // Assignment and access operators
   multi_ptr &operator=(const multi_ptr &) = default;
@@ -441,8 +450,7 @@ class multi_ptr<const void, Space, DecorateAddress> {
   multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(
-            detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}
+      : multi_ptr(Accessor.template get_multi_ptr<DecorateAddress>()) {}
 
   // Only if Space == local_space
   template <
@@ -463,7 +471,7 @@ class multi_ptr<const void, Space, DecorateAddress> {
       typename = typename std::enable_if_t<
           RelaySpace == Space && Space == access::address_space::local_space>>
   multi_ptr(local_accessor<ElementType, Dimensions> Accessor)
-      : m_Pointer(detail::cast_AS<decorated_type *>(Accessor.get_pointer())) {}
+      : multi_ptr(Accessor.template get_multi_ptr<DecorateAddress>()) {}
 
   // Assignment operators
   multi_ptr &operator=(const multi_ptr &) = default;
@@ -567,8 +575,7 @@ class multi_ptr<void, Space, DecorateAddress> {
   multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(
-            detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}
+      : multi_ptr(Accessor.template get_multi_ptr<DecorateAddress>()) {}
 
   // Only if Space == local_space
   template <
@@ -589,7 +596,7 @@ class multi_ptr<void, Space, DecorateAddress> {
       typename = typename std::enable_if_t<
           RelaySpace == Space && Space == access::address_space::local_space>>
   multi_ptr(local_accessor<ElementType, Dimensions> Accessor)
-      : m_Pointer(detail::cast_AS<decorated_type *>(Accessor.get_pointer())) {}
+      : multi_ptr(Accessor.template get_multi_ptr<DecorateAddress>()) {}
 
   // Assignment operators
   multi_ptr &operator=(const multi_ptr &) = default;
@@ -760,7 +767,7 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
                      isPlaceholder, PropertyListT>
                 Accessor) {
-    m_Pointer = detail::cast_AS<pointer_t>(Accessor.get_pointer().get());
+    m_Pointer = detail::cast_AS<pointer_t>(Accessor.get_pointer());
   }
 
   // Only if Space == local_space || generic_space

sycl/test-e2e/Basic/multi_ptr.hpp

@@ -31,7 +31,7 @@ template <typename T> struct point {
 };
 
 template <typename T, access::decorated IsDecorated>
-void innerFunc(id<1> wiID, global_ptr<T, IsDecorated> ptr_1,
+void innerFunc(id<1> wiID, global_ptr<const T, IsDecorated> ptr_1,
                global_ptr<T, IsDecorated> ptr_2,
                global_ptr<T, IsDecorated> ptr_3,
                global_ptr<T, IsDecorated> ptr_4,
@@ -110,9 +110,8 @@ template <typename T, access::decorated IsDecorated> void testMultPtr() {
           private_data[i] = 0;
         localAccessor[wiID.get_local_id()] = 0;
 
-        auto ptr_1 =
-            multi_ptr<T, access::address_space::global_space, IsDecorated>(
-                accessorData_1);
+        auto ptr_1 = multi_ptr<const T, access::address_space::global_space,
+                               IsDecorated>(accessorData_1);
         auto ptr_2 =
             multi_ptr<T, access::address_space::global_space, IsDecorated>(
                 accessorData_2);
@@ -136,19 +135,21 @@ template <typename T, access::decorated IsDecorated> void testMultPtr() {
 
         // Construct extension pointer from accessors.
         auto dev_ptr =
-            multi_ptr<T, access::address_space::ext_intel_global_device_space,
+            multi_ptr<const T,
+                      access::address_space::ext_intel_global_device_space,
                       IsDecorated>(accessorData_1);
-        static_assert(std::is_same_v<ext::intel::device_ptr<T, IsDecorated>,
-                                     decltype(dev_ptr)>,
-                      "Incorrect type for dev_ptr.");
+        static_assert(
+            std::is_same_v<ext::intel::device_ptr<const T, IsDecorated>,
+                           decltype(dev_ptr)>,
+            "Incorrect type for dev_ptr.");
 
         // General conversions in multi_ptr class
         T *RawPtr = nullptr;
         global_ptr<T, IsDecorated> ptr_6 =
             address_space_cast<access::address_space::global_space,
                                IsDecorated>(RawPtr);
 
-        global_ptr<T, IsDecorated> ptr_7(accessorData_1);
+        global_ptr<const T, IsDecorated> ptr_7(accessorData_1);
 
         global_ptr<void, IsDecorated> ptr_8 =
             address_space_cast<access::address_space::global_space,
@@ -206,12 +207,12 @@ void testMultPtrArrowOperator() {
             point<T> private_val = 0;
 
             auto ptr_1 =
-                multi_ptr<point<T>, access::address_space::global_space,
+                multi_ptr<const point<T>, access::address_space::global_space,
                           IsDecorated>(accessorData_1);
             auto ptr_2 = multi_ptr<point<T>, access::address_space::local_space,
                                    IsDecorated>(accessorData_2);
             auto ptr_3 =
-                multi_ptr<point<T>,
+                multi_ptr<const point<T>,
                           access::address_space::ext_intel_global_device_space,
                           IsDecorated>(accessorData_3);
             auto ptr_4 =

sycl/test-e2e/Basic/multi_ptr_legacy.hpp

@@ -8,7 +8,7 @@
 
 #include <cassert>
 #include <iostream>
-#include <sycl/sycl.hpp>
+#include <sycl.hpp>
 #include <type_traits>
 
 using namespace sycl;
@@ -30,7 +30,7 @@ template <typename T> struct point {
 };
 
 template <typename T>
-void innerFunc(id<1> wiID, global_ptr<T> ptr_1, global_ptr<T> ptr_2,
+void innerFunc(id<1> wiID, global_ptr<const T> ptr_1, global_ptr<T> ptr_2,
                local_ptr<T> local_ptr) {
   T t = ptr_1[wiID.get(0)];
   local_ptr[wiID.get(0)] = t;
@@ -64,31 +64,33 @@ template <typename T> void testMultPtr() {
 
       cgh.parallel_for<class testMultPtrKernel<T>>(
           nd_range<1>{10, 10}, [=](nd_item<1> wiID) {
-            auto ptr_1 = make_ptr<T, access::address_space::global_space,
+            auto ptr_1 = make_ptr<const T, access::address_space::global_space,
                                   access::decorated::legacy>(
-                accessorData_1.get_pointer());
+                accessorData_1
+                    .template get_multi_ptr<sycl::access::decorated::legacy>());
             auto ptr_2 = make_ptr<T, access::address_space::global_space,
                                   access::decorated::legacy>(
-                accessorData_2.get_pointer());
+                accessorData_2
+                    .template get_multi_ptr<sycl::access::decorated::legacy>());
             auto local_ptr = make_ptr<T, access::address_space::local_space,
                                       access::decorated::legacy>(
                 localAccessor.get_pointer());
 
             // Construct extension pointer from accessors.
             auto dev_ptr =
-                multi_ptr<T,
+                multi_ptr<const T,
                           access::address_space::ext_intel_global_device_space>(
                     accessorData_1);
-            static_assert(
-                std::is_same_v<ext::intel::device_ptr<T>, decltype(dev_ptr)>,
-                "Incorrect type for dev_ptr.");
+            static_assert(std::is_same_v<ext::intel::device_ptr<const T>,
+                                         decltype(dev_ptr)>,
+                          "Incorrect type for dev_ptr.");
 
             // General conversions in multi_ptr class
             T *RawPtr = nullptr;
             global_ptr<T> ptr_4(RawPtr);
             ptr_4 = RawPtr;
 
-            global_ptr<T> ptr_5(accessorData_1);
+            global_ptr<const T> ptr_5(accessorData_1);
 
             global_ptr<void> ptr_6((void *)RawPtr);
 
@@ -144,9 +146,11 @@ template <typename T> void testMultPtrArrowOperator() {
 
       cgh.parallel_for<class testMultPtrArrowOperatorKernel<T>>(
           sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1>) {
-            auto ptr_1 = make_ptr<point<T>, access::address_space::global_space,
-                                  access::decorated::legacy>(
-                accessorData_1.get_pointer());
+            auto ptr_1 =
+                make_ptr<const point<T>, access::address_space::global_space,
+                         access::decorated::legacy>(
+                    accessorData_1.template get_multi_ptr<
+                        sycl::access::decorated::legacy>());
             auto ptr_2 =
                 make_ptr<point<T>, access::address_space::constant_space,
                          access::decorated::legacy>(
@@ -155,7 +159,7 @@ template <typename T> void testMultPtrArrowOperator() {
                                   access::decorated::legacy>(
                 accessorData_3.get_pointer());
             auto ptr_4 =
-                make_ptr<point<T>,
+                make_ptr<const point<T>,
                          access::address_space::ext_intel_global_device_space,
                          access::decorated::legacy>(
                     accessorData_4.get_pointer());

sycl/test-e2e/GroupAlgorithm/SYCL2020/all_of.cpp

@@ -33,7 +33,9 @@ void test(queue q, InputContainer input, OutputContainer output,
         int lid = it.get_local_id(0);
         out[0] = all_of_group(g, pred(in[lid]));
         out[1] = all_of_group(g, in[lid], pred);
-        out[2] = joint_all_of(g, in.get_pointer(), in.get_pointer() + N, pred);
+        out[2] = joint_all_of(
+            g, in.template get_multi_ptr<access::decorated::no>(),
+            in.template get_multi_ptr<access::decorated::no>() + N, pred);
       });
     });
   }

sycl/test-e2e/GroupAlgorithm/SYCL2020/any_of.cpp

@@ -41,7 +41,9 @@ void test(queue q, InputContainer input, OutputContainer output,
         int lid = it.get_local_id(0);
         out[0] = any_of_group(g, pred(in[lid]));
         out[1] = any_of_group(g, in[lid], pred);
-        out[2] = joint_any_of(g, in.get_pointer(), in.get_pointer() + N, pred);
+        out[2] = joint_any_of(
+            g, in.template get_multi_ptr<access::decorated::no>(),
+            in.template get_multi_ptr<access::decorated::no>() + N, pred);
       });
     });
   }

sycl/test-e2e/GroupAlgorithm/SYCL2020/exclusive_scan.cpp

@@ -92,8 +92,10 @@ void test(queue q, InputContainer input, OutputContainer output,
       accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
       cgh.parallel_for<kernel_name2>(nd_range<1>(G, G), [=](nd_item<1> it) {
         group<1> g = it.get_group();
-        joint_exclusive_scan(g, in.get_pointer(), in.get_pointer() + N,
-                             out.get_pointer(), binary_op);
+        joint_exclusive_scan(
+            g, in.template get_multi_ptr<access::decorated::no>(),
+            in.template get_multi_ptr<access::decorated::no>() + N,
+            out.template get_multi_ptr<access::decorated::no>(), binary_op);
       });
     });
   }
@@ -109,8 +111,11 @@ void test(queue q, InputContainer input, OutputContainer output,
       accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
       cgh.parallel_for<kernel_name3>(nd_range<1>(G, G), [=](nd_item<1> it) {
         group<1> g = it.get_group();
-        joint_exclusive_scan(g, in.get_pointer(), in.get_pointer() + N,
-                             out.get_pointer(), init, binary_op);
+        joint_exclusive_scan(
+            g, in.template get_multi_ptr<access::decorated::no>(),
+            in.template get_multi_ptr<access::decorated::no>() + N,
+            out.template get_multi_ptr<access::decorated::no>(), init,
+            binary_op);
       });
     });
   }