DPC++ Runtime
Runtime libraries for oneAPI DPC++
multi_ptr.hpp
Go to the documentation of this file.
1 //==------------ multi_ptr.hpp - SYCL multi_ptr class ----------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 #include <CL/__spirv/spirv_ops.hpp>
11 #include <cassert>
12 #include <cstddef>
13 #include <sycl/access/access.hpp>
14 #include <sycl/detail/common.hpp>
16 
17 namespace sycl {
19 
20 namespace detail {
21 
22 // Helper to avoid instantiations of invalid non-legacy multi_ptr types.
23 template <typename ElementType, access::address_space Space>
25  using pointer_t =
27  using const_pointer_t = typename multi_ptr<const ElementType, Space,
28  access::decorated::yes>::pointer;
29 };
30 
31 // Specialization for constant_space to avoid creating a non-legacy multi_ptr
32 // with the unsupported space.
33 template <typename ElementType>
34 struct LegacyPointerTypes<ElementType, access::address_space::constant_space> {
36  ElementType, access::address_space::constant_space>::type;
39 };
40 
41 // Helper to avoid instantiations of invalid non-legacy multi_ptr types.
42 template <typename ElementType, access::address_space Space>
44  using reference_t =
47  typename multi_ptr<const ElementType, Space,
48  access::decorated::yes>::reference;
49 };
50 
51 // Specialization for constant_space to avoid creating a non-legacy multi_ptr
52 // with the unsupported space.
53 template <typename ElementType>
54 struct LegacyReferenceTypes<ElementType,
55  access::address_space::constant_space> {
57  ElementType, access::address_space::constant_space>::type;
60 };
61 } // namespace detail
62 
63 // Forward declarations
64 template <typename dataT, int dimensions, access::mode accessMode,
65  access::target accessTarget, access::placeholder isPlaceholder,
66  typename PropertyListT>
67 class accessor;
68 template <typename dataT, int dimensions> class local_accessor;
69 
74 // TODO: Default value for DecorateAddress is for backwards compatiblity. It
75 // should be removed.
76 template <typename ElementType, access::address_space Space,
77  access::decorated DecorateAddress = access::decorated::legacy>
78 class multi_ptr {
79 private:
80  using decorated_type =
82 
83 public:
84  static constexpr bool is_decorated =
85  DecorateAddress == access::decorated::yes;
87 
88  using value_type = ElementType;
89  using pointer = std::conditional_t<is_decorated, decorated_type *,
90  std::add_pointer_t<value_type>>;
91  using reference = std::conditional_t<is_decorated, decorated_type &,
92  std::add_lvalue_reference_t<value_type>>;
93  using iterator_category = std::random_access_iterator_tag;
94  using difference_type = std::ptrdiff_t;
95 
96  static_assert(std::is_same<remove_decoration_t<pointer>,
97  std::add_pointer_t<value_type>>::value);
98  static_assert(std::is_same<remove_decoration_t<reference>,
99  std::add_lvalue_reference_t<value_type>>::value);
100  // Legacy has a different interface.
101  static_assert(DecorateAddress != access::decorated::legacy);
102  // constant_space is only supported in legacy multi_ptr.
103  static_assert(Space != access::address_space::constant_space,
104  "SYCL 2020 multi_ptr does not support the deprecated "
105  "constant_space address space.");
106 
107  // Constructors
108  multi_ptr() : m_Pointer(nullptr) {}
109  multi_ptr(const multi_ptr &) = default;
110  multi_ptr(multi_ptr &&) = default;
111  explicit multi_ptr(typename multi_ptr<ElementType, Space,
112  access::decorated::yes>::pointer ptr)
113  : m_Pointer(ptr) {}
114  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
115 
116  // Only if Space is in
117  // {global_space, ext_intel_global_device_space, generic_space}
118  template <
119  int Dimensions, access::mode Mode, access::placeholder isPlaceholder,
120  typename PropertyListT, access::address_space RelaySpace = Space,
121  typename = typename detail::enable_if_t<
122  RelaySpace == Space &&
123  (Space == access::address_space::generic_space ||
124  Space == access::address_space::global_space ||
125  Space == access::address_space::ext_intel_global_device_space)>>
126  multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
127  isPlaceholder, PropertyListT>
128  Accessor)
129  : multi_ptr(Accessor.get_pointer().get()) {}
130 
131  // Only if Space == local_space || generic_space
132  template <int Dimensions, access::mode Mode,
133  access::placeholder isPlaceholder, typename PropertyListT,
134  access::address_space RelaySpace = Space,
135  typename = typename detail::enable_if_t<
136  RelaySpace == Space &&
137  (Space == access::address_space::generic_space ||
138  Space == access::address_space::local_space)>>
139  multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::local,
140  isPlaceholder, PropertyListT>
141  Accessor)
142  : multi_ptr(Accessor.get_pointer().get()) {}
143 
144  // Only if Space == local_space || generic_space
145  template <int Dimensions, access::address_space RelaySpace = Space,
146  typename = typename detail::enable_if_t<
147  RelaySpace == Space &&
148  (Space == access::address_space::generic_space ||
149  Space == access::address_space::local_space)>>
151  : multi_ptr(Accessor.get_pointer().get()) {}
152 
153  // The following constructors are necessary to create multi_ptr<const
154  // ElementType, Space, DecorateAddress> from accessor<ElementType, ...>.
155  // Constructors above could not be used for this purpose because it will
156  // require 2 implicit conversions of user types which is not allowed by C++:
157  // 1. from accessor<ElementType, ...> to
158  // multi_ptr<ElementType, Space, DecorateAddress>
159  // 2. from multi_ptr<ElementType, Space, DecorateAddress> to
160  // multi_ptr<const ElementType, Space, DecorateAddress>
161 
162  // Only if Space is in
163  // {global_space, ext_intel_global_device_space, generic_space} and element
164  // type is const
165  template <
166  int Dimensions, access::mode Mode, access::placeholder isPlaceholder,
167  typename PropertyListT, access::address_space _Space = Space,
168  typename RelayElementType = ElementType,
169  typename = typename detail::enable_if_t<
170  _Space == Space &&
171  (Space == access::address_space::generic_space ||
172  Space == access::address_space::global_space ||
173  Space == access::address_space::ext_intel_global_device_space) &&
174  std::is_const<RelayElementType>::value &&
175  std::is_same<RelayElementType, ElementType>::value>>
178  Mode, access::target::device, isPlaceholder, PropertyListT>
179  Accessor)
180  : multi_ptr(Accessor.get_pointer().get()) {}
181 
182  // Only if Space == local_space || generic_space and element type is const
183  template <int Dimensions, access::mode Mode,
184  access::placeholder isPlaceholder, typename PropertyListT,
185  access::address_space RelaySpace = Space,
186  typename RelayElementType = ElementType,
187  typename = typename detail::enable_if_t<
188  RelaySpace == Space &&
189  (Space == access::address_space::generic_space ||
190  Space == access::address_space::local_space) &&
191  std::is_const<RelayElementType>::value &&
192  std::is_same<RelayElementType, ElementType>::value>>
195  Mode, access::target::local, isPlaceholder, PropertyListT>
196  Accessor)
197  : multi_ptr(Accessor.get_pointer().get()) {}
198 
199  // Only if Space == local_space || generic_space and element type is const
200  template <int Dimensions, access::address_space RelaySpace = Space,
201  typename RelayElementType = ElementType,
202  typename = typename detail::enable_if_t<
203  RelaySpace == Space &&
204  (Space == access::address_space::generic_space ||
205  Space == access::address_space::local_space) &&
206  std::is_const<RelayElementType>::value &&
207  std::is_same<RelayElementType, ElementType>::value>>
209  Dimensions>
210  Accessor)
211  : multi_ptr(Accessor.get_pointer().get()) {}
212 
213  // Assignment and access operators
214  multi_ptr &operator=(const multi_ptr &) = default;
215  multi_ptr &operator=(multi_ptr &&) = default;
216  multi_ptr &operator=(std::nullptr_t) {
217  m_Pointer = nullptr;
218  return *this;
219  }
220  template <
221  access::address_space OtherSpace, access::decorated OtherIsDecorated,
222  typename =
223  std::enable_if_t<Space == access::address_space::generic_space &&
224  OtherSpace != access::address_space::constant_space>>
225  multi_ptr &
227  m_Pointer = detail::cast_AS<decorated_type *>(Other.get_decorated());
228  return *this;
229  }
230  template <
231  access::address_space OtherSpace, access::decorated OtherIsDecorated,
232  typename =
233  std::enable_if_t<Space == access::address_space::generic_space &&
234  OtherSpace != access::address_space::constant_space>>
235  multi_ptr &
237  m_Pointer = detail::cast_AS<decorated_type *>(std::move(Other.m_Pointer));
238  return *this;
239  }
240 
241  reference operator*() const { return *m_Pointer; }
242  pointer operator->() const { return get(); }
243  reference operator[](difference_type index) const { return m_Pointer[index]; }
244 
245  pointer get() const { return detail::cast_AS<pointer>(m_Pointer); }
246  decorated_type *get_decorated() const { return m_Pointer; }
247  std::add_pointer_t<value_type> get_raw() const {
248  return reinterpret_cast<std::add_pointer_t<value_type>>(get());
249  }
250 
251  __SYCL2020_DEPRECATED("Conversion to pointer type is deprecated since SYCL "
252  "2020. Please use get() instead.")
253  operator pointer() const { return get(); }
254 
255  template <access::address_space OtherSpace,
256  access::decorated OtherIsDecorated,
257  access::address_space RelaySpace = Space,
258  typename = typename std::enable_if_t<
259  RelaySpace == Space &&
260  RelaySpace == access::address_space::generic_space &&
261  (OtherSpace == access::address_space::private_space ||
262  OtherSpace == access::address_space::global_space ||
263  OtherSpace == access::address_space::local_space)>>
266  detail::cast_AS<typename multi_ptr<value_type, OtherSpace,
267  access::decorated::yes>::pointer>(
268  get_decorated())};
269  }
270 
271  template <
272  access::address_space OtherSpace, access::decorated OtherIsDecorated,
273  typename RelayElementType = ElementType,
274  access::address_space RelaySpace = Space,
275  typename = typename std::enable_if_t<
276  std::is_same<RelayElementType, ElementType>::value &&
277  !std::is_const<RelayElementType>::value && RelaySpace == Space &&
278  RelaySpace == access::address_space::generic_space &&
279  (OtherSpace == access::address_space::private_space ||
280  OtherSpace == access::address_space::global_space ||
281  OtherSpace == access::address_space::local_space)>>
282  explicit
285  detail::cast_AS<typename multi_ptr<const value_type, OtherSpace,
286  access::decorated::yes>::pointer>(
287  get_decorated())};
288  }
289 
290  template <access::decorated ConvIsDecorated,
291  typename RelayElementType = ElementType,
292  typename = typename std::enable_if_t<
293  std::is_same<RelayElementType, ElementType>::value &&
294  !std::is_const<RelayElementType>::value>>
298  get_decorated())};
299  }
300 
301  template <access::decorated ConvIsDecorated,
302  typename RelayElementType = ElementType,
303  typename = typename std::enable_if_t<
304  std::is_same<RelayElementType, ElementType>::value &&
305  std::is_const<RelayElementType>::value>>
309  get_decorated())};
310  }
311 
312  template <access::decorated ConvIsDecorated>
315  detail::cast_AS<typename multi_ptr<const value_type, Space,
316  access::decorated::yes>::pointer>(
317  get_decorated())};
318  }
319 
322  return multi_ptr<value_type, Space,
324  get_decorated()};
325  }
326 
327  // Explicit conversion to global_space
328  // Only available if Space == address_space::ext_intel_global_device_space ||
329  // Space == address_space::ext_intel_global_host_space
330  template <
331  access::address_space GlobalSpace = access::address_space::global_space,
332  access::address_space RelaySpace = Space,
333  typename = typename detail::enable_if_t<
334  RelaySpace == Space &&
335  GlobalSpace == access::address_space::global_space &&
336  (Space == access::address_space::ext_intel_global_device_space ||
337  Space == access::address_space::ext_intel_global_host_space)>>
338  explicit
340  using global_pointer_t =
341  typename multi_ptr<ElementType, GlobalSpace,
342  access::decorated::yes>::pointer;
344  detail::cast_AS<global_pointer_t>(get_decorated()));
345  }
346 
347  // Only if Space == global_space
348  template <
349  access::address_space _Space = Space,
350  typename = typename detail::enable_if_t<
351  _Space == Space && Space == access::address_space::global_space>>
352  void prefetch(size_t NumElements) const {
353  size_t NumBytes = NumElements * sizeof(ElementType);
354  using ptr_t = typename detail::DecoratedType<char, Space>::type const *;
355  __spirv_ocl_prefetch(reinterpret_cast<ptr_t>(get_decorated()), NumBytes);
356  }
357 
358  // Arithmetic operators
360  m_Pointer += (difference_type)1;
361  return *this;
362  }
364  multi_ptr result(*this);
365  ++(*this);
366  return result;
367  }
369  m_Pointer -= (difference_type)1;
370  return *this;
371  }
373  multi_ptr result(*this);
374  --(*this);
375  return result;
376  }
378  m_Pointer += r;
379  return *this;
380  }
382  m_Pointer -= r;
383  return *this;
384  }
386  return multi_ptr(get_decorated() + r);
387  }
389  return multi_ptr(get_decorated() - r);
390  }
391 
392 private:
393  decorated_type *m_Pointer;
394 };
395 
397 template <access::address_space Space, access::decorated DecorateAddress>
398 class multi_ptr<const void, Space, DecorateAddress> {
399 private:
400  using decorated_type =
402 
403 public:
404  static constexpr bool is_decorated =
405  DecorateAddress == access::decorated::yes;
407 
408  using value_type = const void;
409  using pointer = std::conditional_t<is_decorated, decorated_type *,
410  std::add_pointer_t<value_type>>;
411  using difference_type = std::ptrdiff_t;
412 
413  static_assert(std::is_same<remove_decoration_t<pointer>,
414  std::add_pointer_t<value_type>>::value);
415  // Legacy has a different interface.
416  static_assert(DecorateAddress != access::decorated::legacy);
417  // constant_space is only supported in legacy multi_ptr.
418  static_assert(Space != access::address_space::constant_space,
419  "SYCL 2020 multi_ptr does not support the deprecated "
420  "constant_space address space.");
421 
422  // Constructors
423  multi_ptr() : m_Pointer(nullptr) {}
424  multi_ptr(const multi_ptr &) = default;
425  multi_ptr(multi_ptr &&) = default;
426  explicit multi_ptr(typename multi_ptr<const void, Space,
427  access::decorated::yes>::pointer ptr)
428  : m_Pointer(ptr) {}
429  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
430 
431  // Only if Space is in
432  // {global_space, ext_intel_global_device_space}
433  template <
434  typename ElementType, int Dimensions, access::mode Mode,
435  access::placeholder isPlaceholder, typename PropertyListT,
436  access::address_space RelaySpace = Space,
437  typename = typename detail::enable_if_t<
438  RelaySpace == Space &&
439  (Space == access::address_space::global_space ||
440  Space == access::address_space::ext_intel_global_device_space)>>
441  multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
442  isPlaceholder, PropertyListT>
443  Accessor)
444  : multi_ptr(Accessor.get_pointer().get()) {}
445 
446  // Only if Space == local_space
447  template <
448  typename ElementType, int Dimensions, access::mode Mode,
449  access::placeholder isPlaceholder, typename PropertyListT,
450  access::address_space RelaySpace = Space,
451  typename = typename detail::enable_if_t<
452  RelaySpace == Space && Space == access::address_space::local_space>>
453  multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::local,
454  isPlaceholder, PropertyListT>
455  Accessor)
456  : multi_ptr(Accessor.get_pointer().get()) {}
457 
458  // Only if Space == local_space
459  template <
460  typename ElementType, int Dimensions,
461  access::address_space RelaySpace = Space,
462  typename = typename detail::enable_if_t<
463  RelaySpace == Space && Space == access::address_space::local_space>>
465  : multi_ptr(Accessor.get_pointer().get()) {}
466 
467  // Assignment operators
468  multi_ptr &operator=(const multi_ptr &) = default;
469  multi_ptr &operator=(multi_ptr &&) = default;
470  multi_ptr &operator=(std::nullptr_t) {
471  m_Pointer = nullptr;
472  return *this;
473  }
474 
475  pointer get() const { return detail::cast_AS<pointer>(m_Pointer); }
476 
477  // Conversion to the underlying pointer type
478  operator pointer() const { return get(); }
479 
480  // Explicit conversion to a multi_ptr<ElementType>
481  template <typename ElementType, typename = typename detail::enable_if_t<
482  std::is_const<ElementType>::value>>
485  detail::cast_AS<typename multi_ptr<ElementType, Space,
486  access::decorated::yes>::pointer>(
487  m_Pointer)};
488  }
489 
490  // Implicit conversion to the negated decoration version of multi_ptr.
493  return multi_ptr<value_type, Space,
495  m_Pointer};
496  }
497 
498  // Explicit conversion to global_space
499  // Only available if Space == address_space::ext_intel_global_device_space ||
500  // Space == address_space::ext_intel_global_host_space
501  template <
502  access::address_space GlobalSpace = access::address_space::global_space,
503  access::address_space RelaySpace = Space,
504  typename = typename detail::enable_if_t<
505  RelaySpace == Space &&
506  GlobalSpace == access::address_space::global_space &&
507  (Space == access::address_space::ext_intel_global_device_space ||
508  Space == access::address_space::ext_intel_global_host_space)>>
509  explicit
511  using global_pointer_t =
512  typename multi_ptr<const void, GlobalSpace,
513  access::decorated::yes>::pointer;
515  detail::cast_AS<global_pointer_t>(m_Pointer));
516  }
517 
518 private:
519  decorated_type *m_Pointer;
520 };
521 
522 // Specialization of multi_ptr for void.
523 template <access::address_space Space, access::decorated DecorateAddress>
525 private:
526  using decorated_type = typename detail::DecoratedType<void, Space>::type;
527 
528 public:
529  static constexpr bool is_decorated =
530  DecorateAddress == access::decorated::yes;
532 
533  using value_type = void;
534  using pointer = std::conditional_t<is_decorated, decorated_type *,
535  std::add_pointer_t<value_type>>;
536  using difference_type = std::ptrdiff_t;
537 
538  static_assert(std::is_same<remove_decoration_t<pointer>,
539  std::add_pointer_t<value_type>>::value);
540  // Legacy has a different interface.
541  static_assert(DecorateAddress != access::decorated::legacy);
542  // constant_space is only supported in legacy multi_ptr.
543  static_assert(Space != access::address_space::constant_space,
544  "SYCL 2020 multi_ptr does not support the deprecated "
545  "constant_space address space.");
546 
547  // Constructors
548  multi_ptr() : m_Pointer(nullptr) {}
549  multi_ptr(const multi_ptr &) = default;
550  multi_ptr(multi_ptr &&) = default;
551  explicit multi_ptr(
553  : m_Pointer(ptr) {}
554  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
555 
556  // Only if Space is in
557  // {global_space, ext_intel_global_device_space}
558  template <
559  typename ElementType, int Dimensions, access::mode Mode,
560  access::placeholder isPlaceholder, typename PropertyListT,
561  access::address_space RelaySpace = Space,
562  typename = typename detail::enable_if_t<
563  RelaySpace == Space &&
564  (Space == access::address_space::global_space ||
565  Space == access::address_space::ext_intel_global_device_space)>>
566  multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
567  isPlaceholder, PropertyListT>
568  Accessor)
569  : multi_ptr(Accessor.get_pointer().get()) {}
570 
571  // Only if Space == local_space
572  template <
573  typename ElementType, int Dimensions, access::mode Mode,
574  access::placeholder isPlaceholder, typename PropertyListT,
575  access::address_space RelaySpace = Space,
576  typename = typename detail::enable_if_t<
577  RelaySpace == Space && Space == access::address_space::local_space>>
578  multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::local,
579  isPlaceholder, PropertyListT>
580  Accessor)
581  : multi_ptr(Accessor.get_pointer().get()) {}
582 
583  // Only if Space == local_space
584  template <
585  typename ElementType, int Dimensions,
586  access::address_space RelaySpace = Space,
587  typename = typename detail::enable_if_t<
588  RelaySpace == Space && Space == access::address_space::local_space>>
590  : multi_ptr(Accessor.get_pointer().get()) {}
591 
592  // Assignment operators
593  multi_ptr &operator=(const multi_ptr &) = default;
594  multi_ptr &operator=(multi_ptr &&) = default;
595  multi_ptr &operator=(std::nullptr_t) {
596  m_Pointer = nullptr;
597  return *this;
598  }
599 
600  pointer get() const { return detail::cast_AS<pointer>(m_Pointer); }
601 
602  // Conversion to the underlying pointer type
603  operator pointer() const { return get(); }
604 
605  // Explicit conversion to a multi_ptr<ElementType>
606  template <typename ElementType>
609  detail::cast_AS<typename multi_ptr<ElementType, Space,
610  access::decorated::yes>::pointer>(
611  m_Pointer)};
612  }
613 
614  // Implicit conversion to the negated decoration version of multi_ptr.
617  return multi_ptr<value_type, Space,
619  m_Pointer};
620  }
621 
622  // Explicit conversion to global_space
623  // Only available if Space == address_space::ext_intel_global_device_space ||
624  // Space == address_space::ext_intel_global_host_space
625  template <
626  access::address_space GlobalSpace = access::address_space::global_space,
627  access::address_space RelaySpace = Space,
628  typename = typename detail::enable_if_t<
629  RelaySpace == Space &&
630  GlobalSpace == access::address_space::global_space &&
631  (Space == access::address_space::ext_intel_global_device_space ||
632  Space == access::address_space::ext_intel_global_host_space)>>
634  using global_pointer_t =
637  detail::cast_AS<global_pointer_t>(m_Pointer));
638  }
639 
640 private:
641  decorated_type *m_Pointer;
642 };
643 
644 // Legacy specialization of multi_ptr.
645 // TODO: Add deprecation warning here when possible.
646 template <typename ElementType, access::address_space Space>
648 public:
649  using element_type =
652  ElementType>;
653  using difference_type = std::ptrdiff_t;
654 
655  // Implementation defined pointer and reference types that correspond to
656  // SYCL/OpenCL interoperability types for OpenCL C functions
657  using pointer_t =
661  using reference_t =
664  typename detail::LegacyReferenceTypes<ElementType,
666 
668 
669  // Constructors
670  multi_ptr() : m_Pointer(nullptr) {}
671  multi_ptr(const multi_ptr &rhs) = default;
672  multi_ptr(multi_ptr &&) = default;
673 #ifdef __SYCL_DEVICE_ONLY__
674  // The generic address space have no corresponding 'opencl_...' attribute and
675  // this constructor is considered as a duplicate for the
676  // multi_ptr(ElementType *pointer) one, so the check is required.
677  template <
678  access::address_space _Space = Space,
679  typename = typename detail::enable_if_t<
680  _Space == Space && Space != access::address_space::generic_space>>
681  multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
682 #endif
683 
684  multi_ptr(ElementType *pointer)
685  : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {
686  // TODO An implementation should reject an argument if the deduced
687  // address space is not compatible with Space.
688  }
689 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
690  template <typename = typename detail::const_if_const_AS<Space, ElementType>>
691  multi_ptr(const ElementType *pointer)
692  : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {}
693 #endif
694 
695  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
696  ~multi_ptr() = default;
697 
698  // Assignment and access operators
699  multi_ptr &operator=(const multi_ptr &) = default;
700  multi_ptr &operator=(multi_ptr &&) = default;
701 
702 #ifdef __SYCL_DEVICE_ONLY__
703  // The generic address space have no corresponding 'opencl_...' attribute and
704  // this operator is considered as a duplicate for the
705  // multi_ptr &operator=(ElementType *pointer) one, so the check is required.
706  template <
707  access::address_space _Space = Space,
708  typename = typename detail::enable_if_t<
709  _Space == Space && Space != access::address_space::generic_space>>
710  multi_ptr &operator=(pointer_t pointer) {
711  m_Pointer = pointer;
712  return *this;
713  }
714 #endif
715 
716  multi_ptr &operator=(ElementType *pointer) {
717  // TODO An implementation should reject an argument if the deduced
718  // address space is not compatible with Space.
719  m_Pointer = detail::cast_AS<pointer_t>(pointer);
720  return *this;
721  }
722 
723  multi_ptr &operator=(std::nullptr_t) {
724  m_Pointer = nullptr;
725  return *this;
726  }
727 
730  using ReturnConstRef = const ElementType &;
731 
733  return *reinterpret_cast<ReturnPtr>(m_Pointer);
734  }
735 
737  return reinterpret_cast<ReturnPtr>(m_Pointer);
738  }
739 
741  return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
742  }
743 
745  return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
746  }
747 
748  // Only if Space is in
749  // {global_space, ext_intel_global_device_space, generic_space}
750  template <
751  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
752  typename PropertyListT, access::address_space _Space = Space,
753  typename = typename detail::enable_if_t<
754  _Space == Space &&
755  (Space == access::address_space::generic_space ||
756  Space == access::address_space::global_space ||
757  Space == access::address_space::ext_intel_global_device_space)>>
758  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
759  isPlaceholder, PropertyListT>
760  Accessor) {
761  m_Pointer = detail::cast_AS<pointer_t>(Accessor.get_pointer().get());
762  }
763 
764  // Only if Space == local_space || generic_space
765  template <
766  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
767  typename PropertyListT, access::address_space _Space = Space,
768  typename = typename detail::enable_if_t<
769  _Space == Space && (Space == access::address_space::generic_space ||
770  Space == access::address_space::local_space)>>
771  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
772  isPlaceholder, PropertyListT>
773  Accessor)
774  : multi_ptr(Accessor.get_pointer()) {}
775 
776  // Only if Space == local_space || generic_space
777  template <int dimensions>
779  : multi_ptr(Accessor.get_pointer()) {}
780 
781  // Only if Space == constant_space
782  template <
783  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
784  typename PropertyListT, access::address_space _Space = Space,
785  typename = typename detail::enable_if_t<
786  _Space == Space && Space == access::address_space::constant_space>>
788  accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
789  isPlaceholder, PropertyListT>
790  Accessor)
791  : multi_ptr(Accessor.get_pointer()) {}
792 
793  // The following constructors are necessary to create multi_ptr<const
794  // ElementType, Space, access::decorated::legacy> from
795  // accessor<ElementType, ...>. Constructors above could not be used for this
796  // purpose because it will require 2 implicit conversions of user types which
797  // is not allowed by C++:
798  // 1. from accessor<ElementType, ...> to
799  // multi_ptr<ElementType, Space, access::decorated::legacy>
800  // 2. from multi_ptr<ElementType, Space, access::decorated::legacy> to
801  // multi_ptr<const ElementType, Space, access::decorated::legacy>
802 
803  // Only if Space is in
804  // {global_space, ext_intel_global_device_space, generic_space} and element
805  // type is const
806  template <
807  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
808  typename PropertyListT, access::address_space _Space = Space,
809  typename ET = ElementType,
810  typename = typename detail::enable_if_t<
811  _Space == Space &&
812  (Space == access::address_space::generic_space ||
813  Space == access::address_space::global_space ||
814  Space == access::address_space::ext_intel_global_device_space) &&
815  std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
816  multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
817  access::target::device, isPlaceholder, PropertyListT>
818  Accessor)
819  : multi_ptr(Accessor.get_pointer()) {}
820 
821  // Only if Space == local_space || generic_space and element type is const
822  template <
823  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
824  typename PropertyListT, access::address_space _Space = Space,
825  typename ET = ElementType,
826  typename = typename detail::enable_if_t<
827  _Space == Space &&
828  (Space == access::address_space::generic_space ||
829  Space == access::address_space::local_space) &&
830  std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
831  multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
832  access::target::local, isPlaceholder, PropertyListT>
833  Accessor)
834  : multi_ptr(Accessor.get_pointer()) {}
835 
836  // Only if Space == local_space || generic_space and element type is const
837  template <
838  int dimensions, access::address_space _Space = Space,
839  typename ET = ElementType,
840  typename = typename detail::enable_if_t<
841  _Space == Space &&
842  (Space == access::address_space::generic_space ||
843  Space == access::address_space::local_space) &&
844  std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
846  local_accessor<typename detail::remove_const_t<ET>, dimensions> Accessor)
847  : multi_ptr(Accessor.get_pointer()) {}
848 
849  // Only if Space == constant_space and element type is const
850  template <
851  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
852  typename PropertyListT, access::address_space _Space = Space,
853  typename ET = ElementType,
854  typename = typename detail::enable_if_t<
855  _Space == Space && Space == access::address_space::constant_space &&
856  std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
858  accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
859  access::target::constant_buffer, isPlaceholder, PropertyListT>
860  Accessor)
861  : multi_ptr(Accessor.get_pointer()) {}
862 
863  // TODO: This constructor is the temporary solution for the existing problem
864  // with conversions from multi_ptr<ElementType, Space,
865  // access::decorated::legacy> to multi_ptr<const ElementType, Space,
866  // access::decorated::legacy>. Without it the compiler fails due to having 3
867  // different same rank paths available.
868  // Constructs multi_ptr<const ElementType, Space, access::decorated::legacy>:
869  // multi_ptr<ElementType, Space, access::decorated::legacy> ->
870  // multi_ptr<const ElementTYpe, Space, access::decorated::legacy>
871  template <typename ET = ElementType>
873  std::is_const<ET>::value && std::is_same<ET, ElementType>::value,
875  access::decorated::legacy>> &ETP)
876  : m_Pointer(ETP.get()) {}
877 
878  // Returns the underlying OpenCL C pointer
879  pointer_t get() const { return m_Pointer; }
880 
881  // Implicit conversion to the underlying pointer type
882  operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); }
883 
884  // Implicit conversion to a multi_ptr<void>
885  // Only available when ElementType is not const-qualified
886  template <typename ET = ElementType>
887  operator multi_ptr<
889  !std::is_const<ET>::value,
890  void>::type,
891  Space, access::decorated::legacy>() const {
892  using ptr_t = typename detail::DecoratedType<void, Space> *;
894  reinterpret_cast<ptr_t>(m_Pointer));
895  }
896 
897  // Implicit conversion to a multi_ptr<const void>
898  // Only available when ElementType is const-qualified
899  template <typename ET = ElementType>
900  operator multi_ptr<
902  std::is_const<ET>::value,
903  const void>::type,
904  Space, access::decorated::legacy>() const {
905  using ptr_t = typename detail::DecoratedType<const void, Space> *;
907  reinterpret_cast<ptr_t>(m_Pointer));
908  }
909 
910  // Implicit conversion to multi_ptr<const ElementType, Space,
911  // access::decorated::legacy>
913  const {
914  using ptr_t =
917  reinterpret_cast<ptr_t>(m_Pointer));
918  }
919 
920  // Arithmetic operators
922  m_Pointer += (difference_type)1;
923  return *this;
924  }
926  multi_ptr result(*this);
927  ++(*this);
928  return result;
929  }
931  m_Pointer -= (difference_type)1;
932  return *this;
933  }
935  multi_ptr result(*this);
936  --(*this);
937  return result;
938  }
940  m_Pointer += r;
941  return *this;
942  }
944  m_Pointer -= r;
945  return *this;
946  }
948  return multi_ptr(m_Pointer + r);
949  }
951  return multi_ptr(m_Pointer - r);
952  }
953 
954 #ifdef __ENABLE_USM_ADDR_SPACE__
955  // Explicit conversion to global_space
956  // Only available if Space == address_space::ext_intel_global_device_space ||
957  // Space == address_space::ext_intel_global_host_space
958  template <
959  access::address_space _Space = Space,
960  typename = typename detail::enable_if_t<
961  _Space == Space &&
962  (Space == access::address_space::ext_intel_global_device_space ||
963  Space == access::address_space::ext_intel_global_host_space)>>
964  explicit operator multi_ptr<ElementType, access::address_space::global_space,
965  access::decorated::legacy>() const {
966  using global_pointer_t = typename detail::DecoratedType<
967  ElementType, access::address_space::global_space>::type *;
968  return multi_ptr<ElementType, access::address_space::global_space,
969  access::decorated::legacy>(
970  reinterpret_cast<global_pointer_t>(m_Pointer));
971  }
972 #endif // __ENABLE_USM_ADDR_SPACE__
973 
974  // Only if Space == global_space
975  template <
976  access::address_space _Space = Space,
977  typename = typename detail::enable_if_t<
978  _Space == Space && Space == access::address_space::global_space>>
979  void prefetch(size_t NumElements) const {
980  size_t NumBytes = NumElements * sizeof(ElementType);
981  using ptr_t = typename detail::DecoratedType<char, Space>::type const *;
982  __spirv_ocl_prefetch(reinterpret_cast<ptr_t>(m_Pointer), NumBytes);
983  }
984 
985 private:
986  pointer_t m_Pointer;
987 };
988 
989 // Legacy specialization of multi_ptr for void.
990 // TODO: Add deprecation warning here when possible.
991 template <access::address_space Space>
992 class multi_ptr<void, Space, access::decorated::legacy> {
993 public:
994  using element_type = void;
995  using difference_type = std::ptrdiff_t;
996 
997  // Implementation defined pointer types that correspond to
998  // SYCL/OpenCL interoperability types for OpenCL C functions
1002 
1004 
1005  // Constructors
1006  multi_ptr() : m_Pointer(nullptr) {}
1007  multi_ptr(const multi_ptr &) = default;
1008  multi_ptr(multi_ptr &&) = default;
1009  multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
1010 #ifdef __SYCL_DEVICE_ONLY__
1011  multi_ptr(void *pointer) : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {
1012  // TODO An implementation should reject an argument if the deduced
1013  // address space is not compatible with Space.
1014  }
1015 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
1016  template <typename = typename detail::const_if_const_AS<Space, void>>
1017  multi_ptr(const void *pointer)
1018  : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {}
1019 #endif
1020 #endif
1021  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
1022  ~multi_ptr() = default;
1023 
1024  // TODO: This constructor is the temporary solution for the existing problem
1025  // with conversions from multi_ptr<ElementType, Space,
1026  // access::decorated::legacy> to multi_ptr<void, Space,
1027  // access::decorated::legacy>. Without it the compiler fails due to having 3
1028  // different same rank paths available.
1029  template <typename ElementType>
1031  : m_Pointer(ETP.get()) {}
1032 
1033  // Assignment operators
1034  multi_ptr &operator=(const multi_ptr &) = default;
1037  m_Pointer = pointer;
1038  return *this;
1039  }
1040 #ifdef __SYCL_DEVICE_ONLY__
1041  multi_ptr &operator=(void *pointer) {
1042  // TODO An implementation should reject an argument if the deduced
1043  // address space is not compatible with Space.
1044  m_Pointer = detail::cast_AS<pointer_t>(pointer);
1045  return *this;
1046  }
1047 #endif
1048  multi_ptr &operator=(std::nullptr_t) {
1049  m_Pointer = nullptr;
1050  return *this;
1051  }
1052 
1053  // Only if Space is in
1054  // {global_space, ext_intel_global_device_space, generic_space}
1055  template <
1056  typename ElementType, int dimensions, access::mode Mode,
1057  typename PropertyListT, access::address_space _Space = Space,
1058  typename = typename detail::enable_if_t<
1059  _Space == Space &&
1060  (Space == access::address_space::generic_space ||
1061  Space == access::address_space::global_space ||
1062  Space == access::address_space::ext_intel_global_device_space)>>
1063  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
1064  access::placeholder::false_t, PropertyListT>
1065  Accessor)
1066  : multi_ptr(Accessor.get_pointer()) {}
1067 
1068  // Only if Space == local_space || generic_space
1069  template <
1070  typename ElementType, int dimensions, access::mode Mode,
1071  typename PropertyListT, access::address_space _Space = Space,
1072  typename = typename detail::enable_if_t<
1073  _Space == Space && (Space == access::address_space::generic_space ||
1074  Space == access::address_space::local_space)>>
1075  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
1076  access::placeholder::false_t, PropertyListT>
1077  Accessor)
1078  : multi_ptr(Accessor.get_pointer()) {}
1079 
1080  // Only if Space == local_space || generic_space
1081  template <
1082  typename ElementType, int dimensions,
1083  access::address_space _Space = Space,
1084  typename = typename detail::enable_if_t<
1085  _Space == Space && (Space == access::address_space::generic_space ||
1086  Space == access::address_space::local_space)>>
1088  : multi_ptr(Accessor.get_pointer()) {}
1089 
1090  // Only if Space == constant_space
1091  template <
1092  typename ElementType, int dimensions, access::mode Mode,
1093  typename PropertyListT, access::address_space _Space = Space,
1094  typename = typename detail::enable_if_t<
1095  _Space == Space && Space == access::address_space::constant_space>>
1097  accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
1098  access::placeholder::false_t, PropertyListT>
1099  Accessor)
1100  : multi_ptr(Accessor.get_pointer()) {}
1101 
1103  // Returns the underlying OpenCL C pointer
1104  pointer_t get() const { return m_Pointer; }
1105 
1106  // Implicit conversion to the underlying pointer type
1107  operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); };
1108 
1109  // Explicit conversion to a multi_ptr<ElementType>
1110  template <typename ElementType>
1111  explicit
1113  using elem_pointer_t =
1116  static_cast<elem_pointer_t>(m_Pointer));
1117  }
1118 
1119  // Implicit conversion to multi_ptr<const void, Space>
1121  using ptr_t = typename detail::DecoratedType<const void, Space>::type *;
1123  reinterpret_cast<ptr_t>(m_Pointer));
1124  }
1125 
1126 private:
1127  pointer_t m_Pointer;
1128 };
1129 
1130 // Legacy specialization of multi_ptr for const void.
1131 // TODO: Add deprecation warning here when possible.
1132 template <access::address_space Space>
1133 class multi_ptr<const void, Space, access::decorated::legacy> {
1134 public:
1135  using element_type = const void;
1136  using difference_type = std::ptrdiff_t;
1137 
1138  // Implementation defined pointer types that correspond to
1139  // SYCL/OpenCL interoperability types for OpenCL C functions
1140  using pointer_t =
1144 
1146 
1147  // Constructors
1148  multi_ptr() : m_Pointer(nullptr) {}
1149  multi_ptr(const multi_ptr &) = default;
1150  multi_ptr(multi_ptr &&) = default;
1151  multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
1152 #ifdef __SYCL_DEVICE_ONLY__
1153  multi_ptr(const void *pointer)
1154  : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {
1155  // TODO An implementation should reject an argument if the deduced
1156  // address space is not compatible with Space.
1157  }
1158 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
1159  template <typename = typename detail::const_if_const_AS<Space, void>>
1160  multi_ptr(const void *pointer)
1161  : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {}
1162 #endif
1163 #endif
1164  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
1165  ~multi_ptr() = default;
1166 
1167  // TODO: This constructor is the temporary solution for the existing problem
1168  // with conversions from multi_ptr<ElementType, Space,
1169  // access::decorated::legacy> to multi_ptr<const void, Space,
1170  // access::decorated::legacy>. Without it the compiler fails due to having 3
1171  // different same rank paths available.
1172  template <typename ElementType>
1174  : m_Pointer(ETP.get()) {}
1175 
1176  // Assignment operators
1177  multi_ptr &operator=(const multi_ptr &) = default;
1180  m_Pointer = pointer;
1181  return *this;
1182  }
1183 #ifdef __SYCL_DEVICE_ONLY__
1184  multi_ptr &operator=(const void *pointer) {
1185  // TODO An implementation should reject an argument if the deduced
1186  // address space is not compatible with Space.
1187  m_Pointer = detail::cast_AS<pointer_t>(pointer);
1188  return *this;
1189  }
1190 #endif
1191  multi_ptr &operator=(std::nullptr_t) {
1192  m_Pointer = nullptr;
1193  return *this;
1194  }
1195 
1196  // Only if Space is in
1197  // {global_space, ext_intel_global_device_space, generic_space}
1198  template <
1199  typename ElementType, int dimensions, access::mode Mode,
1200  typename PropertyListT, access::address_space _Space = Space,
1201  typename = typename detail::enable_if_t<
1202  _Space == Space &&
1203  (Space == access::address_space::generic_space ||
1204  Space == access::address_space::global_space ||
1205  Space == access::address_space::ext_intel_global_device_space)>>
1206  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
1207  access::placeholder::false_t, PropertyListT>
1208  Accessor)
1209  : multi_ptr(Accessor.get_pointer()) {}
1210 
1211  // Only if Space == local_space || generic_space
1212  template <
1213  typename ElementType, int dimensions, access::mode Mode,
1214  typename PropertyListT, access::address_space _Space = Space,
1215  typename = typename detail::enable_if_t<
1216  _Space == Space && (Space == access::address_space::generic_space ||
1217  Space == access::address_space::local_space)>>
1218  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
1219  access::placeholder::false_t, PropertyListT>
1220  Accessor)
1221  : multi_ptr(Accessor.get_pointer()) {}
1222 
1223  // Only if Space == local_space || generic_space
1224  template <
1225  typename ElementType, int dimensions,
1226  access::address_space _Space = Space,
1227  typename = typename detail::enable_if_t<
1228  _Space == Space && (Space == access::address_space::generic_space ||
1229  Space == access::address_space::local_space)>>
1231  : multi_ptr(Accessor.get_pointer()) {}
1232 
1233  // Only if Space == constant_space
1234  template <
1235  typename ElementType, int dimensions, access::mode Mode,
1236  typename PropertyListT, access::address_space _Space = Space,
1237  typename = typename detail::enable_if_t<
1238  _Space == Space && Space == access::address_space::constant_space>>
1240  accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
1241  access::placeholder::false_t, PropertyListT>
1242  Accessor)
1243  : multi_ptr(Accessor.get_pointer()) {}
1244 
1245  // Returns the underlying OpenCL C pointer
1246  pointer_t get() const { return m_Pointer; }
1247 
1248  // Implicit conversion to the underlying pointer type
1249  operator const void *() const {
1250  return reinterpret_cast<const void *>(m_Pointer);
1251  };
1252 
1253  // Explicit conversion to a multi_ptr<const ElementType>
1254  // multi_ptr<const void, Space, access::decorated::legacy> ->
1255  // multi_ptr<const void, Space, access::decorated::legacy>
1256  // The result type must have const specifier.
1257  template <typename ElementType>
1258  explicit
1260  const {
1261  using elem_pointer_t =
1264  static_cast<elem_pointer_t>(m_Pointer));
1265  }
1266 
1267 private:
1268  pointer_t m_Pointer;
1269 };
1270 
1271 #ifdef __cpp_deduction_guides
1272 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
1273  typename PropertyListT, class T>
1274 multi_ptr(accessor<T, dimensions, Mode, access::target::device, isPlaceholder,
1275  PropertyListT>)
1276  -> multi_ptr<T, access::address_space::global_space, access::decorated::no>;
1277 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
1278  typename PropertyListT, class T>
1279 multi_ptr(accessor<T, dimensions, Mode, access::target::constant_buffer,
1280  isPlaceholder, PropertyListT>)
1281  -> multi_ptr<T, access::address_space::constant_space,
1282  access::decorated::legacy>;
1283 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
1284  typename PropertyListT, class T>
1285 multi_ptr(accessor<T, dimensions, Mode, access::target::local, isPlaceholder,
1286  PropertyListT>)
1287  -> multi_ptr<T, access::address_space::local_space, access::decorated::no>;
1288 template <int dimensions, class T>
1289 multi_ptr(local_accessor<T, dimensions>)
1290  -> multi_ptr<T, access::address_space::local_space, access::decorated::no>;
1291 #endif
1292 
1293 template <access::address_space Space, access::decorated DecorateAddress,
1294  typename ElementType>
1295 multi_ptr<ElementType, Space, DecorateAddress>
1296 address_space_cast(ElementType *pointer) {
1297  // TODO An implementation should reject an argument if the deduced address
1298  // space is not compatible with Space.
1299  // Use LegacyPointerTypes here to also allow constant_space
1303  pointer));
1304 }
1305 
1306 template <
1307  typename ElementType, access::address_space Space,
1308  access::decorated DecorateAddress = access::decorated::legacy,
1309  typename = std::enable_if<DecorateAddress == access::decorated::legacy>>
1310 __SYCL2020_DEPRECATED("make_ptr is deprecated since SYCL 2020. Please use "
1311  "address_space_cast instead.")
1312 multi_ptr<ElementType, Space, DecorateAddress> make_ptr(
1313  typename multi_ptr<ElementType, Space, DecorateAddress>::pointer_t
1314  pointer) {
1315  return {pointer};
1316 }
1317 
1318 template <
1319  typename ElementType, access::address_space Space,
1320  access::decorated DecorateAddress,
1321  typename = std::enable_if<DecorateAddress != access::decorated::legacy>>
1322 __SYCL2020_DEPRECATED("make_ptr is deprecated since SYCL 2020. Please use "
1323  "address_space_cast instead.")
1324 multi_ptr<ElementType, Space, DecorateAddress> make_ptr(
1325  typename multi_ptr<ElementType, Space, DecorateAddress>::pointer pointer) {
1326  return address_space_cast<Space, DecorateAddress>(pointer);
1327 }
1328 
1329 #ifdef __SYCL_DEVICE_ONLY__
1330 // An implementation should reject an argument if the deduced address space
1331 // is not compatible with Space.
1332 // This is guaranteed by the c'tor.
1333 template <typename ElementType, access::address_space Space,
1334  access::decorated DecorateAddress = access::decorated::legacy>
1335 __SYCL2020_DEPRECATED("make_ptr is deprecated since SYCL 2020. Please use "
1336  "address_space_cast instead.")
1337 multi_ptr<ElementType, Space, DecorateAddress> make_ptr(ElementType *pointer) {
1338  return address_space_cast<Space, DecorateAddress>(pointer);
1339 }
1340 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
1341 template <typename ElementType, access::address_space Space,
1342  access::decorated DecorateAddress = access::decorated::legacy,
1343  typename = typename detail::const_if_const_AS<Space, ElementType>>
1344 __SYCL2020_DEPRECATED("make_ptr is deprecated since SYCL 2020. Please use "
1345  "address_space_cast instead.")
1346 multi_ptr<ElementType, Space, DecorateAddress> make_ptr(
1347  const ElementType *pointer) {
1348  return multi_ptr<ElementType, Space, DecorateAddress>(pointer);
1349 }
1350 #endif // RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR
1351 #endif // // __SYCL_DEVICE_ONLY__
1352 
1353 template <typename ElementType, access::address_space Space,
1354  access::decorated DecorateAddress>
1357  return lhs.get() == rhs.get();
1358 }
1359 
1360 template <typename ElementType, access::address_space Space,
1361  access::decorated DecorateAddress>
1364  return lhs.get() != rhs.get();
1365 }
1366 
1367 template <typename ElementType, access::address_space Space,
1368  access::decorated DecorateAddress>
1371  return lhs.get() < rhs.get();
1372 }
1373 
1374 template <typename ElementType, access::address_space Space,
1375  access::decorated DecorateAddress>
1378  return lhs.get() > rhs.get();
1379 }
1380 
1381 template <typename ElementType, access::address_space Space,
1382  access::decorated DecorateAddress>
1385  return lhs.get() <= rhs.get();
1386 }
1387 
1388 template <typename ElementType, access::address_space Space,
1389  access::decorated DecorateAddress>
1392  return lhs.get() >= rhs.get();
1393 }
1394 
1395 template <typename ElementType, access::address_space Space,
1396  access::decorated DecorateAddress>
1398  std::nullptr_t) {
1399  return lhs.get() != nullptr;
1400 }
1401 
1402 template <typename ElementType, access::address_space Space,
1403  access::decorated DecorateAddress>
1404 bool operator!=(std::nullptr_t,
1406  return rhs.get() != nullptr;
1407 }
1408 
1409 template <typename ElementType, access::address_space Space,
1410  access::decorated DecorateAddress>
1412  std::nullptr_t) {
1413  return lhs.get() == nullptr;
1414 }
1415 
1416 template <typename ElementType, access::address_space Space,
1417  access::decorated DecorateAddress>
1418 bool operator==(std::nullptr_t,
1420  return rhs.get() == nullptr;
1421 }
1422 
1423 template <typename ElementType, access::address_space Space,
1424  access::decorated DecorateAddress>
1426  std::nullptr_t) {
1427  return lhs.get() != nullptr;
1428 }
1429 
1430 template <typename ElementType, access::address_space Space,
1431  access::decorated DecorateAddress>
1432 bool operator>(std::nullptr_t,
1434  return false;
1435 }
1436 
1437 template <typename ElementType, access::address_space Space,
1438  access::decorated DecorateAddress>
1440  std::nullptr_t) {
1441  return false;
1442 }
1443 
1444 template <typename ElementType, access::address_space Space,
1445  access::decorated DecorateAddress>
1446 bool operator<(std::nullptr_t,
1448  return rhs.get() != nullptr;
1449 }
1450 
1451 template <typename ElementType, access::address_space Space,
1452  access::decorated DecorateAddress>
1454  std::nullptr_t) {
1455  return true;
1456 }
1457 
1458 template <typename ElementType, access::address_space Space,
1459  access::decorated DecorateAddress>
1460 bool operator>=(std::nullptr_t,
1462  return rhs.get() == nullptr;
1463 }
1464 
1465 template <typename ElementType, access::address_space Space,
1466  access::decorated DecorateAddress>
1468  std::nullptr_t) {
1469  return lhs.get() == nullptr;
1470 }
1471 
1472 template <typename ElementType, access::address_space Space,
1473  access::decorated DecorateAddress>
1474 bool operator<=(std::nullptr_t,
1476  return true;
1477 }
1478 
1479 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1480 } // namespace sycl
typename detail::LegacyReferenceTypes< ElementType, Space >::const_reference_t const_reference_t
Definition: multi_ptr.hpp:665
multi_ptr(accessor< typename detail::remove_const_t< ET >, dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:831
ReturnConstRef operator[](difference_type index) const
Definition: multi_ptr.hpp:744
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::device, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:758
detail::const_if_const_AS< Space, ElementType > & ReturnRef
Definition: multi_ptr.hpp:729
typename detail::LegacyPointerTypes< ElementType, Space >::pointer_t pointer_t
Definition: multi_ptr.hpp:658
multi_ptr(accessor< typename detail::remove_const_t< ET >, dimensions, Mode, access::target::constant_buffer, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:857
detail::const_if_const_AS< Space, ElementType > * ReturnPtr
Definition: multi_ptr.hpp:728
multi_ptr(local_accessor< typename detail::remove_const_t< ET >, dimensions > Accessor)
Definition: multi_ptr.hpp:845
multi_ptr(local_accessor< ElementType, dimensions > Accessor)
Definition: multi_ptr.hpp:778
multi_ptr(accessor< typename detail::remove_const_t< ET >, dimensions, Mode, access::target::device, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:816
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::constant_buffer, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:787
typename detail::LegacyPointerTypes< ElementType, Space >::const_pointer_t const_pointer_t
Definition: multi_ptr.hpp:660
multi_ptr(typename detail::enable_if_t< std::is_const< ET >::value &&std::is_same< ET, ElementType >::value, const multi_ptr< typename detail::remove_const_t< ET >, Space, access::decorated::legacy >> &ETP)
Definition: multi_ptr.hpp:872
detail::conditional_t< std::is_same< ElementType, half >::value, sycl::detail::half_impl::BIsRepresentationT, ElementType > element_type
Definition: multi_ptr.hpp:652
typename detail::LegacyReferenceTypes< ElementType, Space >::reference_t reference_t
Definition: multi_ptr.hpp:662
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:771
multi_ptr(accessor< ElementType, Dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:453
multi_ptr(local_accessor< ElementType, Dimensions > Accessor)
Definition: multi_ptr.hpp:464
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
Definition: multi_ptr.hpp:410
multi_ptr & operator=(const multi_ptr &)=default
multi_ptr(accessor< ElementType, Dimensions, Mode, access::target::device, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:441
multi_ptr(typename multi_ptr< const void, Space, access::decorated::yes >::pointer ptr)
Definition: multi_ptr.hpp:426
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::local, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:1218
multi_ptr(const multi_ptr< ElementType, Space, access::decorated::legacy > &ETP)
Definition: multi_ptr.hpp:1173
multi_ptr(local_accessor< ElementType, dimensions > Accessor)
Definition: multi_ptr.hpp:1230
typename detail::LegacyPointerTypes< const void, Space >::pointer_t const_pointer_t
Definition: multi_ptr.hpp:1143
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::constant_buffer, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:1239
typename detail::LegacyPointerTypes< const void, Space >::pointer_t pointer_t
Definition: multi_ptr.hpp:1141
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::device, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:1206
multi_ptr(accessor< ElementType, Dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:578
multi_ptr & operator=(multi_ptr &&)=default
multi_ptr & operator=(const multi_ptr &)=default
multi_ptr(local_accessor< ElementType, Dimensions > Accessor)
Definition: multi_ptr.hpp:589
multi_ptr(accessor< ElementType, Dimensions, Mode, access::target::device, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:566
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
Definition: multi_ptr.hpp:535
multi_ptr(typename multi_ptr< void, Space, access::decorated::yes >::pointer ptr)
Definition: multi_ptr.hpp:551
multi_ptr(local_accessor< ElementType, dimensions > Accessor)
Definition: multi_ptr.hpp:1087
detail::const_if_const_AS< Space, void > * ReturnPtr
Definition: multi_ptr.hpp:1102
typename detail::LegacyPointerTypes< void, Space >::pointer_t pointer_t
Definition: multi_ptr.hpp:999
typename detail::LegacyPointerTypes< const void, Space >::pointer_t const_pointer_t
Definition: multi_ptr.hpp:1001
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::device, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:1063
multi_ptr(const multi_ptr< ElementType, Space, access::decorated::legacy > &ETP)
Definition: multi_ptr.hpp:1030
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::local, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:1075
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::constant_buffer, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:1096
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: multi_ptr.hpp:78
reference operator*() const
Definition: multi_ptr.hpp:241
multi_ptr operator+(difference_type r) const
Definition: multi_ptr.hpp:385
multi_ptr(std::nullptr_t)
Definition: multi_ptr.hpp:114
multi_ptr & operator=(multi_ptr< value_type, OtherSpace, OtherIsDecorated > &&Other)
Definition: multi_ptr.hpp:236
multi_ptr & operator++()
Definition: multi_ptr.hpp:359
reference operator[](difference_type index) const
Definition: multi_ptr.hpp:243
multi_ptr(accessor< ElementType, Dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:139
multi_ptr(const multi_ptr &)=default
__SYCL2020_DEPRECATED("Conversion to pointer type is deprecated since SYCL " "2020. Please use get() instead.") operator pointer() const
Definition: multi_ptr.hpp:251
std::add_pointer_t< value_type > get_raw() const
Definition: multi_ptr.hpp:247
multi_ptr(local_accessor< typename detail::remove_const_t< RelayElementType >, Dimensions > Accessor)
Definition: multi_ptr.hpp:208
multi_ptr(accessor< typename detail::remove_const_t< RelayElementType >, Dimensions, Mode, access::target::device, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:176
multi_ptr & operator=(const multi_ptr &)=default
ElementType value_type
Definition: multi_ptr.hpp:88
multi_ptr operator--(int)
Definition: multi_ptr.hpp:372
multi_ptr(accessor< typename detail::remove_const_t< RelayElementType >, Dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:193
std::ptrdiff_t difference_type
Definition: multi_ptr.hpp:94
multi_ptr operator-(difference_type r) const
Definition: multi_ptr.hpp:388
multi_ptr & operator=(const multi_ptr< value_type, OtherSpace, OtherIsDecorated > &Other)
Definition: multi_ptr.hpp:226
multi_ptr(typename multi_ptr< ElementType, Space, access::decorated::yes >::pointer ptr)
Definition: multi_ptr.hpp:111
multi_ptr(multi_ptr &&)=default
std::random_access_iterator_tag iterator_category
Definition: multi_ptr.hpp:93
void prefetch(size_t NumElements) const
Definition: multi_ptr.hpp:352
multi_ptr operator++(int)
Definition: multi_ptr.hpp:363
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
Definition: multi_ptr.hpp:90
pointer operator->() const
Definition: multi_ptr.hpp:242
multi_ptr & operator-=(difference_type r)
Definition: multi_ptr.hpp:381
multi_ptr & operator+=(difference_type r)
Definition: multi_ptr.hpp:377
multi_ptr & operator=(std::nullptr_t)
Definition: multi_ptr.hpp:216
multi_ptr & operator=(multi_ptr &&)=default
pointer get() const
Definition: multi_ptr.hpp:245
multi_ptr(accessor< ElementType, Dimensions, Mode, access::target::device, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:126
decorated_type * get_decorated() const
Definition: multi_ptr.hpp:246
std::conditional_t< is_decorated, decorated_type &, std::add_lvalue_reference_t< value_type > > reference
Definition: multi_ptr.hpp:92
multi_ptr(local_accessor< ElementType, Dimensions > Accessor)
Definition: multi_ptr.hpp:150
multi_ptr & operator--()
Definition: multi_ptr.hpp:368
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL2020_DEPRECATED(message)
constexpr tuple_element< I, tuple< Types... > >::type & get(sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
ToT cast_AS(FromT from)
Definition: access.hpp:326
typename std::conditional< B, T, F >::type conditional_t
typename std::remove_const< T >::type remove_const_t
typename std::enable_if< B, T >::type enable_if_t
bool operator>=(std::nullptr_t, const multi_ptr< ElementType, Space, DecorateAddress > &rhs)
Definition: multi_ptr.hpp:1460
bool operator<(std::nullptr_t, const multi_ptr< ElementType, Space, DecorateAddress > &rhs)
Definition: multi_ptr.hpp:1446
class __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:2747
bool operator<=(std::nullptr_t, const multi_ptr< ElementType, Space, DecorateAddress > &)
Definition: multi_ptr.hpp:1474
multi_ptr< ElementType, Space, DecorateAddress > address_space_cast(ElementType *pointer)
Definition: multi_ptr.hpp:1296
bool operator!=(std::nullptr_t, const multi_ptr< ElementType, Space, DecorateAddress > &rhs)
Definition: multi_ptr.hpp:1404
DecorateAddress make_ptr(typename multi_ptr< ElementType, Space, DecorateAddress >::pointer pointer)
Definition: multi_ptr.hpp:1324
bool operator>(std::nullptr_t, const multi_ptr< ElementType, Space, DecorateAddress > &)
Definition: multi_ptr.hpp:1432
typename remove_decoration< T >::type remove_decoration_t
Definition: access.hpp:321
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
void __spirv_ocl_prefetch(const char *Ptr, size_t NumBytes) noexcept
Definition: spirv_ops.cpp:47
typename detail::DecoratedType< ElementType, access::address_space::constant_space >::type decorated_type
Definition: multi_ptr.hpp:36
typename multi_ptr< const ElementType, Space, access::decorated::yes >::pointer const_pointer_t
Definition: multi_ptr.hpp:28
typename multi_ptr< ElementType, Space, access::decorated::yes >::pointer pointer_t
Definition: multi_ptr.hpp:26
typename detail::DecoratedType< ElementType, access::address_space::constant_space >::type decorated_type
Definition: multi_ptr.hpp:57
typename multi_ptr< ElementType, Space, access::decorated::yes >::reference reference_t
Definition: multi_ptr.hpp:45
typename multi_ptr< const ElementType, Space, access::decorated::yes >::reference const_reference_t
Definition: multi_ptr.hpp:48
bool operator==(const Slab &Lhs, const Slab &Rhs)