1 /*****************************************************************************
    2 
    3     Copyright (C) Codeplay Software Limited
    4     All Rights Reserved.
    5 
    6     Codeplay's ComputeCpp
    7 
    8 *******************************************************************************/
    9 
   16 #ifndef RUNTIME_INCLUDE_SYCL_MULTI_POINTER_H_
   17 #define RUNTIME_INCLUDE_SYCL_MULTI_POINTER_H_
   18 
   19 #include "SYCL/abacus_types.h"
   20 #include "SYCL/addrspace_cast.h"
   21 #include "SYCL/common.h"
   22 #include "SYCL/deduce.h"
   23 #include "SYCL/type_traits.h"
   24 #include "SYCL/vec.h"
   25 
   26 #include <cstddef>
   27 #include <type_traits>
   28 
   29 namespace cl {
   30 namespace sycl {
   31 
   32 // We need the accessor forward declaration to create pointer class constructor
   33 template <typename elemT, int kDims, access::mode kMode, access::target kTarget,
   34           access::placeholder isPlaceholder>
   35 class accessor;
   36 
   37 namespace detail {
   38 
   45 template <typename dataType, cl::sycl::access::address_space asp>
   46 struct address_space_trait {
   55   template <cl::sycl::access::address_space delay>
   56   struct delayed_assert {
   57     enum { value = false };
   58   };
   59   static_assert(delayed_assert<asp>::value, "Unhandled address space.");
   60 };
   61 
   66 template <typename dataType>
   67 struct address_space_trait<dataType,
   68                            cl::sycl::access::address_space::private_space> {
   69   using original_type = dataType;
   70   using address_space_type = COMPUTECPP_CL_ASP_PRIVATE dataType;
   71 // ocl_qualified_address_space_type is the same as address_space_type (i.e. the
   72 // type with the appropriate asp qualifier) with the exception of the private
   73 // ASP - address_space_type is not qualified, while
   74 // ocl_qualified_address_space_type is (in clang, unqualified != private)
   75 #if defined __SYCL_DEVICE_ONLY__
   76   using ocl_qualified_address_space_type =
   77       __attribute__((opencl_private)) dataType;
   78 #else
   79   using ocl_qualified_address_space_type = address_space_type;
   80 #endif
   81 
   84   static constexpr bool hasTarget = false;
   85   static constexpr auto targetValue = COMPUTECPP_ACCESS_TARGET_DEVICE;
   86 };
   87 
   92 template <typename dataType>
   93 struct address_space_trait<dataType,
   94                            cl::sycl::access::address_space::global_space> {
   95   using original_type = dataType;
   96   using address_space_type = COMPUTECPP_CL_ASP_GLOBAL dataType;
   97   using ocl_qualified_address_space_type = address_space_type;
   98   static constexpr bool hasTarget = true;
   99   static constexpr auto targetValue = COMPUTECPP_ACCESS_TARGET_DEVICE;
  100 };
  101 
  106 template <typename dataType>
  107 struct address_space_trait<dataType,
  108                            cl::sycl::access::address_space::constant_space> {
  109 #if defined(COMPUTECPP_EXT_READ_ACC_CONST_PTR_ENABLE)
  110   using original_type = const dataType;
  111   using address_space_type = const COMPUTECPP_CL_ASP_CONSTANT dataType;
  112 #else
  113   using original_type = dataType;
  114   using address_space_type = COMPUTECPP_CL_ASP_CONSTANT dataType;
  115 #endif  // COMPUTECPP_EXT_READ_ACC_CONST_PTR_ENABLE
  116 
  117   using ocl_qualified_address_space_type = address_space_type;
  118   static constexpr bool hasTarget = true;
  119   COMPUTECPP_HOST_CXX_DIAGNOSTIC(push)
  120   COMPUTECPP_SET_IGNORE_DEPRECATED_API
  121   static constexpr auto targetValue = access::target::constant_buffer;
  122   COMPUTECPP_HOST_CXX_DIAGNOSTIC(pop)
  123 };
  124 
  129 template <typename dataType>
  130 struct address_space_trait<dataType,
  131                            cl::sycl::access::address_space::local_space> {
  132   using original_type = dataType;
  133   using address_space_type = COMPUTECPP_CL_ASP_LOCAL dataType;
  134   using ocl_qualified_address_space_type = address_space_type;
  135   static constexpr bool hasTarget = true;
  136   COMPUTECPP_HOST_CXX_DIAGNOSTIC(push)
  137   COMPUTECPP_SET_IGNORE_DEPRECATED_API
  138   static constexpr auto targetValue = access::target::local;
  139   COMPUTECPP_HOST_CXX_DIAGNOSTIC(pop)
  140 };
  141 
  147 template <typename dataType>
  148 struct address_space_trait<dataType,
  149                            access::address_space::subgroup_local_space> {
  150   using original_type = dataType;
  151   using address_space_type = COMPUTECPP_CL_ASP_SUBGROUP_LOCAL dataType;
  152   using ocl_qualified_address_space_type = address_space_type;
  153   static constexpr bool hasTarget = true;
  154   static constexpr auto targetValue = access::target::subgroup_local;
  155 };
  156 
  165 template <typename dataType, cl::sycl::access::address_space asp,
  166           typename DerivedT>
  167 class vec_ptr_class_base {};
  168 
  178 template <typename dataType, int width, cl::sycl::access::address_space asp,
  179           typename DerivedT>
  180 class vec_ptr_class_base<vec<dataType, width>, asp, DerivedT> {
  181   DerivedT& derived() { return static_cast<DerivedT&>(*this); }
  182 
  183  public:
  184 #ifdef __SYCL_DEVICE_ONLY__
  185   using vector_asp_type = typename detail::address_space_trait<
  186       detail::__sycl_vector<dataType, width>, asp>::address_space_type;
  187 
  188   operator vector_asp_type*() {
  189     return reinterpret_cast<vector_asp_type*>(derived().m_elem);
  190   }
  191 #else
  192 
  199   using abacus_type =
  200       typename abacus::convert_abacus_sycl<vec<dataType, width>>::abacus_type;
  201 
  202   operator abacus_type*() {
  203     abacus_type* abacusVec = reinterpret_cast<abacus_type*>(derived().m_elem);
  204     return abacusVec;
  205   }
  206 #endif
  207 };
  208 
  209 #ifdef __SYCL_DEVICE_ONLY__
  210 
  217 #ifdef __COMPUTECPP_NO_ASP__
  218 template <typename dataType, cl::sycl::access::address_space asp>
  219 using address_space_trait_maybe =
  220     typename std::enable_if<asp == access::address_space::subgroup_local_space,
  221                             address_space_trait<dataType, asp>>::type;
  222 #else   //  __COMPUTECPP_NO_ASP__
  223 template <typename dataType, cl::sycl::access::address_space asp>
  224 using address_space_trait_maybe = address_space_trait<dataType, asp>;
  225 #endif  //  __COMPUTECPP_NO_ASP__
  226 
  227 // The following are functions that take an address space qualified pointer,
  228 // and return a non-qualified one. The function we need to call for each depends
  229 // on how and which compiler is being used
  230 
  231 // OFFLOAD (Both regular, and the NO-ASP workaround)
  232 #ifndef __SYCL_COMPUTECPP_ASP__
  233 template <typename dataType, cl::sycl::access::address_space asp>
  234 typename detail::address_space_trait<dataType, asp>::original_type*
  235 get_pointer_internal_type(
  236     typename detail::address_space_trait<dataType, asp>::address_space_type*
  237         ptr) {
  238   // The body is only used in __COMPUTECPP_NO_ASP__ mode with non-subgroup ptrs
  239   typename detail::address_space_trait<dataType, asp>::original_type* tmp;
  240   *(detail::reinterpret_addrspace_cast<decltype(ptr)*>(&tmp)) = ptr;
  241   return tmp;
  242 }
  243 
  244 template <typename dataType, cl::sycl::access::address_space asp>
  245 typename detail::address_space_trait<
  246     dataType, asp>::address_space_type* __attribute__((__offload__))
  247 get_pointer_internal_type(typename detail::address_space_trait_maybe<
  248                           dataType, asp>::address_space_type* ptr) {
  249   return ptr;
  250 }
  251 
  252 template <typename dataType, cl::sycl::access::address_space asp>
  253 typename detail::address_space_trait<
  254     dataType, asp>::address_space_type* __attribute__((__offload__(2)))
  255 get_pointer_internal_type(typename detail::address_space_trait_maybe<
  256                           dataType, asp>::address_space_type* ptr) {
  257   return ptr;
  258 }
  259 #else   //  __SYCL_COMPUTECPP_ASP__
  260 // ASP
  261 template <typename dataType, cl::sycl::access::address_space asp>
  262 typename detail::address_space_trait<dataType, asp>::original_type*
  263 get_pointer_internal_type(
  264     typename detail::address_space_trait<dataType, asp>::address_space_type*
  265         ptr) {
  266   // ASP can decay address space qualified pointers to unqualified
  267   return ptr;
  268 }
  269 #endif  //  __SYCL_COMPUTECPP_ASP__
  270 #else   //  __SYCL_DEVICE_ONLY__
  271 // HOST
  274 template <typename dataType, cl::sycl::access::address_space asp>
  275 typename detail::address_space_trait<dataType, asp>::address_space_type*
  276 get_pointer_internal_type(
  277     typename detail::address_space_trait<dataType, asp>::address_space_type*
  278         ptr) {
  279   return ptr;
  280 }
  281 #endif  // __SYCL_DEVICE_ONLY__
  282 
  296 // Type helper for hidden types (used when __COMPUTECPP_NO_ASP__ is set)
  297 template <bool hidden>
  298 struct get_visible_type {};
  299 
  300 template <>
  301 struct get_visible_type<true> {
  306   template <typename aspPtr, typename origPtr>
  307   static origPtr get(aspPtr /*asp*/, origPtr orig) {
  308     return orig;
  309   }
  310 };
  311 
  312 template <>
  313 struct get_visible_type<false> {
  314   template <typename aspPtr, typename origPtr>
  315   static aspPtr get(aspPtr asp, origPtr /*orig*/) {
  316     return asp;
  317   }
  318 };
  319 
  321 // multi_ptr_base
  322 
  331 template <typename dataType, cl::sycl::access::address_space asp>
  332 class multi_ptr_base
  333     : public detail::vec_ptr_class_base<
  334           typename address_space_trait<dataType, asp>::original_type, asp,
  335           multi_ptr_base<dataType, asp>> {
  336  private:
  337   using base_t = detail::vec_ptr_class_base<
  338       typename address_space_trait<dataType, asp>::original_type, asp,
  339       multi_ptr_base<dataType, asp>>;
  340 
  341  protected:
  347 #ifdef __COMPUTECPP_NO_ASP__
  348   constexpr static bool hasHiddenAddrSpace =
  349       asp != cl::sycl::access::address_space::subgroup_local_space;
  350 #else
  351   constexpr static bool hasHiddenAddrSpace = false;
  352 #endif
  353 
  356   using element_type =
  357       typename address_space_trait<dataType, asp>::original_type;
  358 
  359   using asp_type =
  360       typename detail::address_space_trait<element_type,
  361                                            asp>::address_space_type;
  362   using original_type_ptr =
  363       typename detail::address_space_trait<element_type, asp>::original_type*;
  364 
  367   using difference_type = std::ptrdiff_t;
  368 
  371   using asp_pointer_t = asp_type*;
  372 
  375   using pointer_t =
  376       typename std::conditional<hasHiddenAddrSpace, original_type_ptr,
  377                                 asp_pointer_t>::type;
  378 
  381   using const_pointer_t = const pointer_t;
  382 
  383   using value_type = element_type;
  384   using pointer = pointer_t;
  385   using iterator_category = std::random_access_iterator_tag;
  386 
  387  public:
  390   static constexpr cl::sycl::access::address_space address_space = asp;
  391 
  393   // Constructors
  394 
  397   multi_ptr_base() : m_elem(nullptr) {}
  398 
  402   multi_ptr_base(pointer_t ptr)
  403       : m_elem(reinterpret_cast<asp_pointer_t>(ptr)) {}
  404 
  405  protected:
  406   friend base_t;
  407 
  408   original_type_ptr get_pointer_internal() const {
  409     return detail::get_pointer_internal_type<element_type, asp>(m_elem);
  410   }
  411 
  412   original_type_ptr get_pointer_internal() const volatile {
  413     return detail::get_pointer_internal_type<element_type, asp>(m_elem);
  414   }
  415 
  419   inline void increment_pointer(difference_type r) { m_elem += r; }
  420 
  421  public:
  423   // Pointer retrieval
  424 
  427   inline pointer_t get() const {
  428     return get_visible_type<hasHiddenAddrSpace>::get(m_elem,
  429                                                      get_pointer_internal());
  430   }
  431 
  434   original_type_ptr operator->() const { return get_pointer_internal(); }
  435 
  437   // Conversion operators
  438 
  441   operator original_type_ptr() const { return get_pointer_internal(); }
  442 
  445   operator original_type_ptr() const volatile { return get_pointer_internal(); }
  446 
  447 #if defined(__SYCL_DEVICE_ONLY__)
  448 
  450   template <typename U = pointer_t>
  451   operator typename std::enable_if<!std::is_same<original_type_ptr, U>::value &&
  452                                        !hasHiddenAddrSpace,
  453                                    U>::type() {
  454     return m_elem;
  455   }
  458   template <typename U = pointer_t>
  459   operator typename std::enable_if<!std::is_same<original_type_ptr, U>::value &&
  460                                        !hasHiddenAddrSpace,
  461                                    U>::type() const {
  462     return m_elem;
  463   }
  466   template <typename U = pointer_t>
  467   operator typename std::enable_if<!std::is_same<original_type_ptr, U>::value &&
  468                                        !hasHiddenAddrSpace,
  469                                    U>::type() volatile {
  470     return m_elem;
  471   }
  474   template <typename U = pointer_t>
  475   operator typename std::enable_if<!std::is_same<original_type_ptr, U>::value &&
  476                                        !hasHiddenAddrSpace,
  477                                    U>::type() const volatile {
  478     return m_elem;
  479   }
  480 #endif  // defined(__SYCL_DEVICE_ONLY__)
  481 
  483   // prefetch
  484 
  492   template <COMPUTECPP_ENABLE_IF(dataType,
  493                                  (asp == access::address_space::global_space))>
  494   void prefetch(size_t /*numElements*/) const {
  495     // Note: not implemented
  496   }
  497 
  498  private:
  501   asp_pointer_t m_elem;
  502 };
  503 
  504 }  // namespace detail
  505 
  507 // multi_ptr
  508 
  517 template <typename dataType, cl::sycl::access::address_space asp>
  518 class multi_ptr : public detail::multi_ptr_base<dataType, asp> {
  519  protected:
  520   using multi_ptr_base = detail::multi_ptr_base<dataType, asp>;
  521 
  522   using original_type_ref =
  523       typename detail::address_space_trait<dataType, asp>::original_type&;
  524   using original_type_cref =
  525       const typename detail::address_space_trait<dataType, asp>::original_type&;
  526 
  529   using non_const_data_t = typename std::remove_const<dataType>::type;
  530 
  531  public:
  533 
  537   using ptr_t = typename multi_ptr_base::pointer;
  538 
  542   using ptr_unqual_t = typename multi_ptr_base::element_type*;
  543 
  545 
  548   using reference_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(202001, "Use reference") =
  549       typename multi_ptr_base::asp_type&;
  550 
  553   using const_reference_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(
  554       202001,
  555       "Use add_const_t<reference>") = const typename multi_ptr_base::asp_type&;
  556 
  557   using element_type COMPUTECPP_DEPRECATED_BY_SYCL_VER(
  558       202001, "Use value_type") = typename multi_ptr_base::element_type;
  559   using difference_type = typename multi_ptr_base::difference_type;
  560   using pointer_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(202001,
  561                                                     "Use pointer") = ptr_t;
  562   using const_pointer_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(
  563       202001, "Use std::add_const_t<pointer>") =
  564       typename multi_ptr_base::const_pointer_t;
  565 
  566 #if SYCL_LANGUAGE_VERSION >= 202001
  567   using value_type = typename multi_ptr_base::value_type;
  570   using pointer = typename multi_ptr_base::pointer;
  572   using reference = typename multi_ptr_base::asp_type&;
  574   using iterator_category = typename multi_ptr_base::iterator_category;
  575 #endif  // SYCL_LANGUAGE_VERSION
  576 
  578   // Constructors
  579 
  582   multi_ptr() = default;
  583 
  587   multi_ptr(ptr_t ptr) : multi_ptr_base(ptr) {}
  588 
  593   template <COMPUTECPP_ENABLE_IF(dataType, (std::is_const<dataType>::value))>
  594   multi_ptr(const multi_ptr<non_const_data_t, asp>& rhs)
  595       : multi_ptr_base(rhs.get()) {}
  596 
  600 #if !defined(__SYCL_COMPUTECPP_ASP__) && defined(__SYCL_DEVICE_ONLY__)
  601   // Offload
  602   multi_ptr(const multi_ptr<void, asp>& rhs)
  603       : multi_ptr_base(static_cast<ptr_unqual_t>(rhs.get())) {}
  604 #else   // Host + ASP
  605   multi_ptr(const multi_ptr<void, asp>& rhs)
  606       : multi_ptr_base(static_cast<ptr_t>(rhs.get())) {}
  607 #endif  // !defined(__SYCL_COMPUTECPP_ASP__) && defined(__SYCL_DEVICE_ONLY__)
  608 
  613 #if !defined(__SYCL_COMPUTECPP_ASP__) && defined(__SYCL_DEVICE_ONLY__)
  614   // Offload
  615   template <COMPUTECPP_ENABLE_IF(dataType, (std::is_const<dataType>::value))>
  616   multi_ptr(const multi_ptr<const void, asp>& rhs)
  617       : multi_ptr_base(static_cast<const ptr_unqual_t>(rhs.get())) {}
  618 #else   // Host + ASP
  619   template <COMPUTECPP_ENABLE_IF(dataType, (std::is_const<dataType>::value))>
  620   multi_ptr(const multi_ptr<const void, asp>& rhs)
  621       : multi_ptr_base(static_cast<const ptr_t>(rhs.get())) {}
  622 #endif  // !defined(__SYCL_COMPUTECPP_ASP__) && defined(__SYCL_DEVICE_ONLY__)
  623 
  635   template <int dimensions, access::mode Mode,
  636             access::placeholder isPlaceholder = access::placeholder::false_t,
  637             COMPUTECPP_ENABLE_IF(dataType, (detail::address_space_trait<
  638                                                dataType, asp>::hasTarget))>
  639   multi_ptr(cl::sycl::accessor<
  640             non_const_data_t, dimensions, Mode,
  641             detail::address_space_trait<non_const_data_t, asp>::targetValue,
  642             isPlaceholder>
  643                 acc)
  644       : multi_ptr(acc.get_pointer()) {}
  645 
  657   template <int dimensions, access::mode Mode,
  658             access::placeholder isPlaceholder = access::placeholder::false_t,
  659             COMPUTECPP_ENABLE_IF(dataType, (detail::address_space_trait<
  660                                                 dataType, asp>::hasTarget &&
  661                                             std::is_const<dataType>::value))>
  662   multi_ptr(
  663       cl::sycl::accessor<
  664           const non_const_data_t, dimensions, Mode,
  665           detail::address_space_trait<const non_const_data_t, asp>::targetValue,
  666           isPlaceholder>
  667           acc)
  668       : multi_ptr(acc.get_pointer()) {}
  669 
  684   template <COMPUTECPP_ENABLE_IF(dataType,
  685                                  (!std::is_same<ptr_t, ptr_unqual_t>::value))>
  686 #ifdef __SYCL_COMPUTECPP_ASP__
  687   multi_ptr(ptr_unqual_t ptr) : multi_ptr(detail::addrspace_cast<ptr_t>(ptr)) {
  688   }
  689 #else   // Offload
  690   multi_ptr(ptr_unqual_t ptr);
  691 #endif  // __SYCL_COMPUTECPP_ASP__
  692 
  694   // Conversions
  695 
  696  protected:
  698   using void_conversion_t =
  699       typename std::conditional<std::is_const<dataType>::value, const void,
  700                                 void>::type;
  701 
  702  public:
  705   operator multi_ptr<void_conversion_t, asp>() const {
  706     return multi_ptr<void_conversion_t, asp>(
  707         reinterpret_cast<void_conversion_t*>(this->get()));
  708   }
  709 
  711   // Reference retrieval
  712 
  715   original_type_ref operator*() { return *this->get_pointer_internal(); }
  716 
  719   original_type_cref operator*() const { return *this->get_pointer_internal(); }
  720 
  725   COMPUTECPP_DEPRECATED_BY_SYCL_VER(
  726       201703, "multi_ptr::operator[] is no longer available.")
  727   original_type_ref operator[](size_t i) {
  728     return this->get_pointer_internal()[i];
  729   }
  730 
  735   COMPUTECPP_DEPRECATED_BY_SYCL_VER(
  736       201703, "multi_ptr::operator[] is no longer available.")
  737   original_type_cref operator[](size_t i) const {
  738     return this->get_pointer_internal()[i];
  739   }
  740 
  742   // Arithmetic operators
  743 
  747   inline multi_ptr& operator++() {
  748     this->increment_pointer(1);
  749     return *this;
  750   }
  751 
  756   inline multi_ptr operator++(int) {
  757     auto copy = multi_ptr(*this);
  758     this->operator++();
  759     return copy;
  760   }
  761 
  765   inline multi_ptr& operator--() {
  766     this->increment_pointer(-1);
  767     return *this;
  768   }
  769 
  774   inline multi_ptr operator--(int) {
  775     auto copy = multi_ptr(*this);
  776     this->operator--();
  777     return copy;
  778   }
  779 
  784   inline multi_ptr& operator+=(difference_type r) {
  785     this->increment_pointer(r);
  786     return *this;
  787   }
  788 
  793   inline multi_ptr operator+(difference_type r) const {
  794     return multi_ptr(this->get_pointer_internal() + r);
  795   }
  796 
  801   inline multi_ptr& operator-=(difference_type r) {
  802     this->increment_pointer(-r);
  803     return *this;
  804   }
  805 
  810   inline multi_ptr operator-(difference_type r) const {
  811     return multi_ptr(this->get_pointer_internal() - r);
  812   }
  813 };
  814 
  815 namespace detail {
  819 #ifdef __SYCL_DEVICE_ONLY__
  820 #ifdef __SYCL_COMPUTECPP_ASP__
  821 template <typename dataT, cl::sycl::access::address_space asp>
  822 typename address_space_trait<dataT, asp>::ocl_qualified_address_space_type*
  823 multi_ptr_get_internal_type(multi_ptr<dataT, asp> mp) {
  824   return detail::addrspace_cast<typename address_space_trait<
  825       dataT, asp>::ocl_qualified_address_space_type*>(mp.get());
  826 }
  827 #else   // Offload
  828 template <typename dataT, cl::sycl::access::address_space asp>
  829 typename address_space_trait<dataT, asp>::ocl_qualified_address_space_type*
  830 multi_ptr_get_internal_type(multi_ptr<dataT, asp> mp) {
  831   return reinterpret_cast<typename address_space_trait<
  832       dataT, asp>::ocl_qualified_address_space_type*>(mp.get());
  833 }
  834 #endif  // __SYCL_COMPUTECPP_ASP__
  835 #else   //  __SYCL_DEVICE_ONLY__
  836 template <typename dataT, cl::sycl::access::address_space asp>
  837 typename multi_ptr<dataT, asp>::pointer_t multi_ptr_get_internal_type(
  838     multi_ptr<dataT, asp> mp) {
  839   return mp.get();
  840 }
  841 #endif  //  __SYCL_DEVICE_ONLY__
  842 }  // namespace detail
  843 
  847 template <cl::sycl::access::address_space asp>
  848 class multi_ptr<void, asp> : public detail::multi_ptr_base<void, asp> {
  849  protected:
  850   using multi_ptr_base = detail::multi_ptr_base<void, asp>;
  851 
  852  public:
  854 
  858   using ptr_t = typename multi_ptr_base::pointer;
  859 
  863   using ptr_unqual_t = typename multi_ptr_base::element_type*;
  864 
  866 
  867   using element_type COMPUTECPP_DEPRECATED_BY_SYCL_VER(
  868       202001, "Use value_type") = typename multi_ptr_base::element_type;
  869   using difference_type = typename multi_ptr_base::difference_type;
  870   using pointer_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(202001, "Use pointer") =
  871       typename multi_ptr_base::pointer_t;
  872   using const_pointer_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(
  873       202001, "Use std::add_const_t<pointer>") =
  874       typename multi_ptr_base::const_pointer_t;
  875 
  876 #if SYCL_LANGUAGE_VERSION >= 202001
  877   using value_type = typename multi_ptr_base::value_type;
  880   using pointer = typename multi_ptr_base::pointer;
  882   using iterator_category = typename multi_ptr_base::iterator_category;
  883 #endif  // SYCL_LANGUAGE_VERSION
  884 
  886   // Constructors
  887 
  890   multi_ptr() = default;
  891 
  895   multi_ptr(ptr_t ptr) : multi_ptr_base(ptr) {}
  896 
  907   template <int dimensions, access::mode Mode,
  908             access::placeholder isPlaceholder = access::placeholder::false_t,
  909             COMPUTECPP_ENABLE_IF(
  910                 void, (detail::address_space_trait<void, asp>::hasTarget))>
  911   multi_ptr(cl::sycl::accessor<
  912             void, dimensions, Mode,
  913             detail::address_space_trait<void, asp>::targetValue, isPlaceholder>
  914                 acc)
  915       : multi_ptr(acc.get_pointer()) {}
  916 
  931   template <COMPUTECPP_ENABLE_IF(void,
  932                                  (!std::is_same<ptr_t, ptr_unqual_t>::value))>
  933 #ifdef __SYCL_COMPUTECPP_ASP__
  934   multi_ptr(ptr_unqual_t ptr) : multi_ptr(detail::addrspace_cast<ptr_t>(ptr)) {
  935   }
  936 #else   // Offload
  937   multi_ptr(ptr_unqual_t ptr);
  938 #endif  // __SYCL_COMPUTECPP_ASP__
  939 
  945   template <typename ElementType,
  946             COMPUTECPP_ENABLE_IF(ElementType,
  947                                  (!std::is_same<ElementType, void>::value))>
  948   explicit multi_ptr(const multi_ptr<ElementType, asp>& ptr)
  949       : multi_ptr_base(static_cast<ptr_t>(ptr.get())) {}
  950 
  962   template <typename ElementType, int dimensions, access::mode Mode,
  963             access::placeholder isPlaceholder,
  964             COMPUTECPP_ENABLE_IF(
  965                 ElementType,
  966                 (detail::address_space_trait<ElementType, asp>::hasTarget))>
  967   multi_ptr(accessor<ElementType, dimensions, Mode,
  968                      detail::address_space_trait<ElementType, asp>::targetValue,
  969                      isPlaceholder>
  970                 acc)
  971       : multi_ptr_base(acc.get_pointer().get()) {}
  972 
  976   template <typename ElementType>
  977   explicit operator multi_ptr<ElementType, asp>() const {
  978     using elem_ptr_t =
  979         typename detail::address_space_trait<ElementType, asp>::original_type*;
  980     return multi_ptr<ElementType, asp>(
  981         static_cast<elem_ptr_t>(this->get_pointer_internal()));
  982   }
  983 };
  984 
  988 template <cl::sycl::access::address_space asp>
  989 class multi_ptr<const void, asp>
  990     : public detail::multi_ptr_base<const void, asp> {
  991  protected:
  992   using multi_ptr_base = detail::multi_ptr_base<const void, asp>;
  993 
  994  public:
  996 
 1000   using ptr_t = typename multi_ptr_base::pointer;
 1001 
 1005   using ptr_unqual_t = typename multi_ptr_base::element_type*;
 1006 
 1008 
 1009   using element_type COMPUTECPP_DEPRECATED_BY_SYCL_VER(
 1010       202001, "Use value_type") = typename multi_ptr_base::element_type;
 1011   using difference_type = typename multi_ptr_base::difference_type;
 1012   using pointer_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(202001, "Use pointer") =
 1013       typename multi_ptr_base::pointer_t;
 1014   using const_pointer_t COMPUTECPP_DEPRECATED_BY_SYCL_VER(
 1015       202001, "Use std::add_const_t<pointer>") =
 1016       typename multi_ptr_base::const_pointer_t;
 1017 
 1018 #if SYCL_LANGUAGE_VERSION >= 202001
 1019   using value_type = typename multi_ptr_base::value_type;
 1022   using pointer = typename multi_ptr_base::pointer;
 1024   using iterator_category = typename multi_ptr_base::iterator_category;
 1025 #endif  // SYCL_LANGUAGE_VERSION
 1026 
 1028   // Constructors
 1029 
 1032   multi_ptr() = default;
 1033 
 1037   multi_ptr(ptr_t ptr) : multi_ptr_base(ptr) {}
 1038 
 1049   template <int dimensions, access::mode Mode,
 1050             access::placeholder isPlaceholder = access::placeholder::false_t,
 1051             COMPUTECPP_ENABLE_IF(const void, (detail::address_space_trait<
 1052                                                  const void, asp>::hasTarget))>
 1053   multi_ptr(cl::sycl::accessor<
 1054             const void, dimensions, Mode,
 1055             detail::address_space_trait<const void, asp>::targetValue,
 1056             isPlaceholder>
 1057                 acc)
 1058       : multi_ptr(acc.get_pointer()) {}
 1059 
 1074   template <COMPUTECPP_ENABLE_IF(const void,
 1075                                  (!std::is_same<ptr_t, ptr_unqual_t>::value))>
 1076 #ifdef __SYCL_COMPUTECPP_ASP__
 1077   multi_ptr(ptr_unqual_t ptr) : multi_ptr(detail::addrspace_cast<ptr_t>(ptr)) {
 1078   }
 1079 #else   // Offload
 1080   multi_ptr(ptr_unqual_t ptr);
 1081 #endif  // __SYCL_COMPUTECPP_ASP__
 1082 
 1088   template <typename ElementType,
 1089             COMPUTECPP_ENABLE_IF(
 1090                 ElementType, (!std::is_same<ElementType, const void>::value))>
 1091   explicit multi_ptr(const multi_ptr<ElementType, asp>& ptr)
 1092       : multi_ptr_base(static_cast<ptr_t>(ptr.get())) {}
 1093 
 1105   template <typename ElementType, int dimensions, access::mode Mode,
 1106             access::placeholder isPlaceholder,
 1107             COMPUTECPP_ENABLE_IF(
 1108                 ElementType,
 1109                 (detail::address_space_trait<ElementType, asp>::hasTarget))>
 1110   multi_ptr(accessor<ElementType, dimensions, Mode,
 1111                      detail::address_space_trait<ElementType, asp>::targetValue,
 1112                      isPlaceholder>
 1113                 acc)
 1114       : multi_ptr_base(acc.get_pointer().get()) {}
 1115 
 1119   template <typename ElementType>
 1120   explicit operator multi_ptr<ElementType, asp>() const {
 1121     using elem_ptr_t =
 1122         typename detail::address_space_trait<ElementType, asp>::original_type*;
 1123     return multi_ptr<ElementType, asp>(
 1124         static_cast<elem_ptr_t>(this->get_pointer_internal()));
 1125   }
 1126 };
 1127 
 1129 // Explicit pointer class aliases
 1130 
 1135 template <typename dataType>
 1136 using global_ptr = multi_ptr<dataType, access::address_space::global_space>;
 1137 
 1142 template <typename dataType>
 1143 using local_ptr = multi_ptr<dataType, access::address_space::local_space>;
 1144 
 1149 template <typename dataType>
 1150 using private_ptr = multi_ptr<dataType, access::address_space::private_space>;
 1151 
 1156 template <typename dataType>
 1157 using constant_ptr = multi_ptr<dataType, access::address_space::constant_space>;
 1158 
 1159 namespace codeplay {
 1164 template <typename dataType>
 1165 using subgroup_local_ptr =
 1166     multi_ptr<dataType, access::address_space::subgroup_local_space>;
 1167 }  // namespace codeplay
 1168 
 1170 // make_ptr
 1171 
 1178 template <typename dataType, cl::sycl::access::address_space Space>
 1179 typename std::enable_if<
 1180     !std::is_same<typename multi_ptr<dataType, Space>::ptr_t, dataType*>::value,
 1181     multi_ptr<dataType, Space>>::type
 1182 make_ptr(dataType* ptr) {
 1183   return multi_ptr<dataType, Space>(ptr);
 1184 }
 1185 
 1186 template <typename dataType, cl::sycl::access::address_space Space>
 1187 multi_ptr<dataType, Space> make_ptr(
 1188     typename multi_ptr<dataType, Space>::ptr_t ptr) {
 1189   return multi_ptr<dataType, Space>(ptr);
 1190 }
 1191 
 1192 namespace detail {
 1193 
 1194 template <typename T, access::address_space AddressSpace>
 1195 multi_ptr<deduce_type_t<T>, AddressSpace> deduce_type_impl_f(
 1196     multi_ptr<T, AddressSpace>);
 1197 
 1200 template <typename T, access::address_space AddressSpace>
 1201 struct deduce_type<multi_ptr<T, AddressSpace>> {
 1202   using type = multi_ptr<deduce_type_t<T>, AddressSpace>;
 1203 };
 1204 
 1207 template <typename T, access::address_space AddressSpace>
 1208 struct deduce_type<multi_ptr<const T, AddressSpace>> {
 1209   using type = multi_ptr<const deduce_type_t<T>, AddressSpace>;
 1210 };
 1211 
 1214 template <typename T, access::address_space AddressSpace>
 1215 struct deduce_type<multi_ptr<volatile T, AddressSpace>> {
 1216   using type = multi_ptr<volatile deduce_type_t<T>, AddressSpace>;
 1217 };
 1218 
 1221 template <typename T, access::address_space AddressSpace>
 1222 struct deduce_type<multi_ptr<const volatile T, AddressSpace>> {
 1223   using type = multi_ptr<const volatile deduce_type_t<T>, AddressSpace>;
 1224 };
 1225 
 1226 }  // namespace detail
 1227 
 1228 }  // namespace sycl
 1229 }  // namespace cl
 1230 
 1231 #if SYCL_LANGUAGE_VERSION >= 202001
 1232 namespace std {
 1233 
 1238 template <typename dataType, cl::sycl::access::address_space asp>
 1239 struct iterator_traits<cl::sycl::multi_ptr<dataType, asp>> {
 1240   using value_type = dataType;
 1241   using pointer = dataType*;
 1242   using reference = dataType&;
 1243   using difference_type = ptrdiff_t;
 1244   using iterator_category =
 1245       typename cl::sycl::multi_ptr<dataType, asp>::iterator_category;
 1246 };
 1247 
 1248 }  // namespace std
 1249 #endif  // SYCL_LANGUAGE_VERSION >= 202001
 1250 
 1251 /******************************************/
 1252 /* Operators definitions: explicit pointer */
 1253 /******************************************/
 1254 
 1256 
 1260 #define COMPUTECPP_EXPLICIT_POINTER_OP_DEFINITION(explicit_ptr)                \
 1261   template <typename dataType>                                                 \
 1262   bool operator==(const explicit_ptr<dataType>& lhs,                           \
 1263                   const explicit_ptr<dataType>& rhs) {                         \
 1264     return (lhs.get() == rhs.get());                                           \
 1265   }                                                                            \
 1266   template <typename dataType>                                                 \
 1267   bool operator!=(const explicit_ptr<dataType>& lhs,                           \
 1268                   const explicit_ptr<dataType>& rhs) {                         \
 1269     return !(lhs == rhs);                                                      \
 1270   }                                                                            \
 1271   template <typename dataType>                                                 \
 1272   bool operator<(const explicit_ptr<dataType>& lhs,                            \
 1273                  const explicit_ptr<dataType>& rhs) {                          \
 1274     return (lhs.get() < rhs.get());                                            \
 1275   }                                                                            \
 1276   template <typename dataType>                                                 \
 1277   bool operator>(const explicit_ptr<dataType>& lhs,                            \
 1278                  const explicit_ptr<dataType>& rhs) {                          \
 1279     return (lhs.get() > rhs.get());                                            \
 1280   }                                                                            \
 1281   template <typename dataType>                                                 \
 1282   bool operator>=(const explicit_ptr<dataType>& lhs,                           \
 1283                   const explicit_ptr<dataType>& rhs) {                         \
 1284     return (lhs == rhs) || (lhs > rhs);                                        \
 1285   }                                                                            \
 1286   template <typename dataType>                                                 \
 1287   bool operator<=(const explicit_ptr<dataType>& lhs,                           \
 1288                   const explicit_ptr<dataType>& rhs) {                         \
 1289     return (lhs == rhs) || (lhs < rhs);                                        \
 1290   }                                                                            \
 1291   template <typename dataType>                                                 \
 1292   bool operator==(const explicit_ptr<dataType>& lhs, std::nullptr_t rhs) {     \
 1293     (void)rhs;                                                                 \
 1294     return (lhs.get() == rhs);                                                 \
 1295   }                                                                            \
 1296   template <typename dataType>                                                 \
 1297   bool operator==(std::nullptr_t lhs, const explicit_ptr<dataType>& rhs) {     \
 1298     (void)lhs;                                                                 \
 1299     return rhs == lhs;                                                         \
 1300   }                                                                            \
 1301   template <typename dataType>                                                 \
 1302   bool operator!=(const explicit_ptr<dataType>& lhs, std::nullptr_t rhs) {     \
 1303     (void)rhs;                                                                 \
 1304     return !(lhs == rhs);                                                      \
 1305   }                                                                            \
 1306   template <typename dataType>                                                 \
 1307   bool operator!=(std::nullptr_t lhs, const explicit_ptr<dataType>& rhs) {     \
 1308     (void)lhs;                                                                 \
 1309     return !(lhs == rhs);                                                      \
 1310   }                                                                            \
 1311   template <typename dataType>                                                 \
 1312   bool operator>(const explicit_ptr<dataType>& lhs, std::nullptr_t rhs) {      \
 1313     (void)rhs;                                                                 \
 1314  /* anything > nullptr => nullptr != anything*/                             \
 1315     return lhs != rhs;                                                         \
 1316   }                                                                            \
 1317   template <typename dataType>                                                 \
 1318   bool operator>(std::nullptr_t, const explicit_ptr<dataType>&) {              \
 1319  /*  nullptr > anything => false */                                         \
 1320     return false;                                                              \
 1321   }                                                                            \
 1322   template <typename dataType>                                                 \
 1323   bool operator<(const explicit_ptr<dataType>&, std::nullptr_t) {              \
 1324  /*  anything < nullptr => false */                                         \
 1325     return false;                                                              \
 1326   }                                                                            \
 1327   template <typename dataType>                                                 \
 1328   bool operator<(std::nullptr_t lhs, const explicit_ptr<dataType>& rhs) {      \
 1329     (void)lhs;                                                                 \
 1330  /*  nullptr < anything => nullptr != anything */                           \
 1331     return lhs != rhs;                                                         \
 1332   }                                                                            \
 1333   template <typename dataType>                                                 \
 1334   bool operator>=(const explicit_ptr<dataType>&, std::nullptr_t) {             \
 1335     return true;                                                               \
 1336   }                                                                            \
 1337   template <typename dataType>                                                 \
 1338   bool operator>=(std::nullptr_t lhs, const explicit_ptr<dataType>& rhs) {     \
 1339     (void)lhs;                                                                 \
 1340  /*  nullptr >= anything => nullptr == anything */                          \
 1341     return lhs == rhs;                                                         \
 1342   }                                                                            \
 1343   template <typename dataType>                                                 \
 1344   bool operator<=(const explicit_ptr<dataType>& lhs, std::nullptr_t rhs) {     \
 1345     (void)rhs;                                                                 \
 1346  /*  anything <= nullptr => anything == nullptr */                          \
 1347     return lhs == rhs;                                                         \
 1348   }                                                                            \
 1349   template <typename dataType>                                                 \
 1350   bool operator<=(std::nullptr_t, const explicit_ptr<dataType>&) {             \
 1351  /*  nullptr <= anything => anything == nullptr */                          \
 1352     return true;                                                               \
 1353   }
 1354 
 1355 COMPUTECPP_EXPLICIT_POINTER_OP_DEFINITION(cl::sycl::private_ptr)
 1356 COMPUTECPP_EXPLICIT_POINTER_OP_DEFINITION(cl::sycl::global_ptr)
 1357 COMPUTECPP_EXPLICIT_POINTER_OP_DEFINITION(cl::sycl::constant_ptr)
 1358 COMPUTECPP_EXPLICIT_POINTER_OP_DEFINITION(cl::sycl::local_ptr)
 1359 COMPUTECPP_EXPLICIT_POINTER_OP_DEFINITION(
 1360     cl::sycl::codeplay::subgroup_local_ptr)
 1361 
 1362 #undef COMPUTECPP_EXPLICIT_POINTER_OP_DEFINITION
 1363 
 1365 
 1366 /***********************************/
 1367 /* Operators definitions: multi_ptr */
 1368 /***********************************/
 1369 
 1374 template <typename dataType, cl::sycl::access::address_space Space>
 1375 bool operator==(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1376                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1377   return lhs.get_pointer() == rhs.get_pointer();
 1378 }
 1379 
 1385 template <typename dataType, cl::sycl::access::address_space Space>
 1386 bool operator!=(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1387                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1388   return lhs.get_pointer() != rhs.get_pointer();
 1389 }
 1390 
 1395 template <typename dataType, cl::sycl::access::address_space Space>
 1396 bool operator<(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1397                const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1398   return lhs.get_pointer() < rhs.get_pointer();
 1399 }
 1400 
 1405 template <typename dataType, cl::sycl::access::address_space Space>
 1406 bool operator>(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1407                const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1408   return lhs.get_pointer() > rhs.get_pointer();
 1409 }
 1410 
 1415 template <typename dataType, cl::sycl::access::address_space Space>
 1416 bool operator<=(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1417                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1418   return lhs.get_pointer() <= rhs.get_pointer();
 1419 }
 1420 
 1425 template <typename dataType, cl::sycl::access::address_space Space>
 1426 bool operator>=(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1427                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1428   return lhs.get_pointer() >= rhs.get_pointer();
 1429 }
 1430 
 1434 template <typename dataType, cl::sycl::access::address_space Space>
 1435 bool operator!=(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1436                 std::nullptr_t rhs) {
 1437   (void)rhs;
 1438   return lhs.get_pointer() != rhs;
 1439 }
 1440 
 1444 template <typename dataType, cl::sycl::access::address_space Space>
 1445 bool operator!=(std::nullptr_t lhs,
 1446                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1447   (void)lhs;
 1448   return lhs != rhs.get_pointer();
 1449 }
 1450 
 1454 template <typename dataType, cl::sycl::access::address_space Space>
 1455 bool operator==(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1456                 std::nullptr_t rhs) {
 1457   (void)rhs;
 1458   return lhs.get_pointer() == rhs;
 1459 }
 1460 
 1464 template <typename dataType, cl::sycl::access::address_space Space>
 1465 bool operator==(std::nullptr_t lhs,
 1466                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1467   (void)lhs;
 1468   return lhs == rhs.get_pointer();
 1469 }
 1470 
 1474 template <typename dataType, cl::sycl::access::address_space Space>
 1475 bool operator>(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1476                std::nullptr_t rhs) {
 1477   (void)rhs;
 1478   return lhs.get_pointer() > rhs;
 1479 }
 1480 
 1484 template <typename dataType, cl::sycl::access::address_space Space>
 1485 bool operator>(std::nullptr_t lhs,
 1486                const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1487   (void)lhs;
 1488   return lhs > rhs.get_pointer();
 1489 }
 1490 
 1494 template <typename dataType, cl::sycl::access::address_space Space>
 1495 bool operator<(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1496                std::nullptr_t rhs) {
 1497   (void)rhs;
 1498   return lhs.get_pointer() < rhs;
 1499 }
 1500 
 1504 template <typename dataType, cl::sycl::access::address_space Space>
 1505 bool operator<(std::nullptr_t lhs,
 1506                const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1507   (void)lhs;
 1508   return lhs < rhs.get_pointer();
 1509 }
 1510 
 1514 template <typename dataType, cl::sycl::access::address_space Space>
 1515 bool operator>=(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1516                 std::nullptr_t rhs) {
 1517   (void)rhs;
 1518   return lhs.get_pointer() >= rhs;
 1519 }
 1520 
 1524 template <typename dataType, cl::sycl::access::address_space Space>
 1525 bool operator>=(std::nullptr_t lhs,
 1526                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1527   (void)lhs;
 1528   return lhs >= rhs.get_pointer();
 1529 }
 1530 
 1534 template <typename dataType, cl::sycl::access::address_space Space>
 1535 bool operator<=(const cl::sycl::multi_ptr<dataType, Space>& lhs,
 1536                 std::nullptr_t rhs) {
 1537   (void)rhs;
 1538   return lhs.get_pointer() <= rhs;
 1539 }
 1540 
 1544 template <typename dataType, cl::sycl::access::address_space Space>
 1545 bool operator<=(std::nullptr_t lhs,
 1546                 const cl::sycl::multi_ptr<dataType, Space>& rhs) {
 1547   (void)lhs;
 1548   return lhs <= rhs.get_pointer();
 1549 }
 1550 
 1551 #endif  // RUNTIME_INCLUDE_SYCL_MULTI_POINTER_H_