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