[SYCL2020] Update the multi_ptr class

include/hipSYCL/sycl/access.hpp

@@ -151,6 +151,12 @@ enum class fence_space : char {
   global_and_local
 };
 
+enum class decorated : char {
+  no,
+  yes,
+  legacy // [[deprecated]] TODO: This is deprecated but used as default value for multi_ptr template argument
+};
+
 inline std::ostream &operator<<(std::ostream &out,
                          const sycl::access::placeholder value)
 {

include/hipSYCL/sycl/libkernel/accessor.hpp

@@ -645,6 +645,21 @@ inline constexpr sycl::access_mode default_access_mode() {
   return std::is_const_v<T> ? access_mode::read : access_mode::read_write;
 }
 
+template <target accessTarget, typename value_type>
+struct multi_ptr_for_target {
+  using type = void*;
+};
+
+template <typename value_type>
+struct multi_ptr_for_target<target::device, value_type> {
+  using type = global_ptr<value_type>;
+};
+
+template <typename value_type>
+struct multi_ptr_for_target<target::local, value_type> {
+  using type = local_ptr<value_type>;
+};
+
 } // detail
 
 namespace property {
@@ -772,7 +787,8 @@ class HIPSYCL_EMPTY_BASES accessor
       typename detail::accessor::accessor_data_type<dataT, accessmode>::value;
   using reference = value_type &;
   using const_reference = const dataT &;
-  // TODO accessor_ptr
+  template <access::decorated IsDecorated>
+  using accessor_ptr = typename detail::multi_ptr_for_target<accessTarget, value_type>::type;
   using iterator = detail::accessor_iterator<value_type, dimensions, accessor>;
   using const_iterator = detail::accessor_iterator<const value_type, dimensions, accessor>;
   using reverse_iterator = std::reverse_iterator<iterator>;
@@ -1178,8 +1194,11 @@ class HIPSYCL_EMPTY_BASES accessor
   }
 
   /* Available only when: accessTarget == access::target::host_buffer */
+  // TODO: T == host_buffer is deprecated in SYCL2020
   template<access::target T = accessTarget,
-           typename = std::enable_if_t<T==access::target::host_buffer>>
+           typename = std::enable_if_t<
+	     T == access::target::host_buffer ||
+	     T == access::target::host_task>>
   dataT *get_pointer() const noexcept
   {
     return const_cast<dataT*>(this->_ptr.get());
@@ -1203,6 +1222,16 @@ class HIPSYCL_EMPTY_BASES accessor
     return constant_ptr<dataT>{const_cast<dataT*>(this->_ptr.get())};
   }
 
+  /* Available only when: accessTarget == access::target::device */
+  template<access::decorated IsDecorated,
+           access::target T = accessTarget,
+           typename = std::enable_if_t<T == access::target::device>>
+  HIPSYCL_UNIVERSAL_TARGET
+  accessor_ptr<IsDecorated> get_multi_ptr() const noexcept
+  {
+    return accessor_ptr<IsDecorated>{this->_ptr.get()};
+  }
+
   iterator begin() const noexcept {
     return iterator::make_begin(this);
   }
@@ -2126,10 +2155,6 @@ class accessor<
   range<dimensions> _num_elements;
 };
 
-template <typename dataT, int dimensions = 1>
-using local_accessor = accessor<dataT, dimensions, access::mode::read_write,
-  access::target::local>;
-
 namespace detail::accessor {
 
 template<class AccessorType>