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