DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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>
14 #include <cassert>
15 #include <cstddef>
16 
18 namespace sycl {
19 // Forward declaration
20 template <typename dataT, int dimensions, access::mode accessMode,
21  access::target accessTarget, access::placeholder isPlaceholder,
22  typename PropertyListT>
23 class accessor;
24 
29 template <typename ElementType, access::address_space Space> class multi_ptr {
30 public:
31  using element_type =
32  detail::conditional_t<std::is_same<ElementType, half>::value,
34  ElementType>;
35  using difference_type = std::ptrdiff_t;
36 
37  // Implementation defined pointer and reference types that correspond to
38  // SYCL/OpenCL interoperability types for OpenCL C functions
40  using const_pointer_t =
42  using reference_t =
44  using const_reference_t =
46 
47  static constexpr access::address_space address_space = Space;
48 
49  // Constructors
50  multi_ptr() : m_Pointer(nullptr) {}
51  multi_ptr(const multi_ptr &rhs) = default;
52  multi_ptr(multi_ptr &&) = default;
53 #ifdef __SYCL_DEVICE_ONLY__
54  multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
55 #endif
56 
57  multi_ptr(ElementType *pointer) : m_Pointer((pointer_t)(pointer)) {
58  // TODO An implementation should reject an argument if the deduced
59  // address space is not compatible with Space.
60  }
61 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
62  template <typename = typename detail::const_if_const_AS<Space, ElementType>>
63  multi_ptr(const ElementType *pointer) : m_Pointer((pointer_t)(pointer)) {}
64 #endif
65 
66  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
67  ~multi_ptr() = default;
68 
69  // Assignment and access operators
70  multi_ptr &operator=(const multi_ptr &) = default;
71  multi_ptr &operator=(multi_ptr &&) = default;
72 
73 #ifdef __SYCL_DEVICE_ONLY__
74  multi_ptr &operator=(pointer_t pointer) {
75  m_Pointer = pointer;
76  return *this;
77  }
78 #endif
79 
80  multi_ptr &operator=(ElementType *pointer) {
81  // TODO An implementation should reject an argument if the deduced
82  // address space is not compatible with Space.
83  m_Pointer = (pointer_t)pointer;
84  return *this;
85  }
86 
87  multi_ptr &operator=(std::nullptr_t) {
88  m_Pointer = nullptr;
89  return *this;
90  }
91 
94  using ReturnConstRef = const ElementType &;
95 
96  ReturnRef operator*() const {
97  return *reinterpret_cast<ReturnPtr>(m_Pointer);
98  }
99 
101  return reinterpret_cast<ReturnPtr>(m_Pointer);
102  }
103 
105  return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
106  }
107 
109  return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
110  }
111 
112  // Only if Space == global_space || global_device_space
113  template <int dimensions, access::mode Mode,
114  access::placeholder isPlaceholder, typename PropertyListT,
115  access::address_space _Space = Space,
116  typename = typename detail::enable_if_t<
117  _Space == Space &&
118  (Space == access::address_space::global_space ||
119  Space == access::address_space::global_device_space)>>
121  accessor<ElementType, dimensions, Mode, access::target::global_buffer,
122  isPlaceholder, PropertyListT>
123  Accessor) {
124  m_Pointer = (pointer_t)(Accessor.get_pointer().get());
125  }
126 
127  // Only if Space == local_space
128  template <int dimensions, access::mode Mode,
129  access::placeholder isPlaceholder, typename PropertyListT,
130  access::address_space _Space = Space,
131  typename = typename detail::enable_if_t<
132  _Space == Space && Space == access::address_space::local_space>>
133  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
134  isPlaceholder, PropertyListT>
135  Accessor)
136  : multi_ptr(Accessor.get_pointer()) {}
137 
138  // Only if Space == constant_space
139  template <
140  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
141  typename PropertyListT, access::address_space _Space = Space,
142  typename = typename detail::enable_if_t<
143  _Space == Space && Space == access::address_space::constant_space>>
145  accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
146  isPlaceholder, PropertyListT>
147  Accessor)
148  : multi_ptr(Accessor.get_pointer()) {}
149 
150  // The following constructors are necessary to create multi_ptr<const
151  // ElementType, Space> from accessor<ElementType, ...>. Constructors above
152  // could not be used for this purpose because it will require 2 implicit
153  // conversions of user types which is not allowed by C++:
154  // 1. from accessor<ElementType, ...> to multi_ptr<ElementType, Space>
155  // 2. from multi_ptr<ElementType, Space> to multi_ptr<const ElementType,
156  // Space>
157 
158  // Only if Space == global_space || global_device_space and element type is
159  // const
160  template <
161  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
162  typename PropertyListT, access::address_space _Space = Space,
163  typename ET = ElementType,
164  typename = typename detail::enable_if_t<
165  _Space == Space &&
166  (Space == access::address_space::global_space ||
167  Space == access::address_space::global_device_space) &&
168  std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
170  accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
171  access::target::global_buffer, isPlaceholder, PropertyListT>
172  Accessor)
173  : multi_ptr(Accessor.get_pointer()) {}
174 
175  // Only if Space == local_space and element type is const
176  template <
177  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
178  typename PropertyListT, access::address_space _Space = Space,
179  typename ET = ElementType,
180  typename = typename detail::enable_if_t<
181  _Space == Space && Space == access::address_space::local_space &&
182  std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
183  multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
184  access::target::local, isPlaceholder, PropertyListT>
185  Accessor)
186  : multi_ptr(Accessor.get_pointer()) {}
187 
188  // Only if Space == constant_space and element type is const
189  template <
190  int dimensions, access::mode Mode, access::placeholder isPlaceholder,
191  typename PropertyListT, access::address_space _Space = Space,
192  typename ET = ElementType,
193  typename = typename detail::enable_if_t<
194  _Space == Space && Space == access::address_space::constant_space &&
195  std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
197  accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
198  access::target::constant_buffer, isPlaceholder, PropertyListT>
199  Accessor)
200  : multi_ptr(Accessor.get_pointer()) {}
201 
202  // TODO: This constructor is the temporary solution for the existing problem
203  // with conversions from multi_ptr<ElementType, Space> to
204  // multi_ptr<const ElementType, Space>. Without it the compiler
205  // fails due to having 3 different same rank paths available.
206  // Constructs multi_ptr<const ElementType, Space>:
207  // multi_ptr<ElementType, Space> -> multi_ptr<const ElementTYpe, Space>
208  template <typename ET = ElementType>
210  std::is_const<ET>::value && std::is_same<ET, ElementType>::value,
211  const multi_ptr<typename detail::remove_const_t<ET>, Space>> &ETP)
212  : m_Pointer(ETP.get()) {}
213 
214  // Returns the underlying OpenCL C pointer
215  pointer_t get() const { return m_Pointer; }
216 
217  // Implicit conversion to the underlying pointer type
218  operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); }
219 
220  // Implicit conversion to a multi_ptr<void>
221  // Only available when ElementType is not const-qualified
222  template <typename ET = ElementType>
223  operator multi_ptr<
225  !std::is_const<ET>::value,
226  void>::type,
227  Space>() const {
228  using ptr_t = typename detail::DecoratedType<void, Space> *;
229  return multi_ptr<void, Space>(reinterpret_cast<ptr_t>(m_Pointer));
230  }
231 
232  // Implicit conversion to a multi_ptr<const void>
233  // Only available when ElementType is const-qualified
234  template <typename ET = ElementType>
235  operator multi_ptr<
237  std::is_const<ET>::value,
238  const void>::type,
239  Space>() const {
240  using ptr_t = typename detail::DecoratedType<const void, Space> *;
241  return multi_ptr<const void, Space>(reinterpret_cast<ptr_t>(m_Pointer));
242  }
243 
244  // Implicit conversion to multi_ptr<const ElementType, Space>
246  using ptr_t =
249  reinterpret_cast<ptr_t>(m_Pointer));
250  }
251 
252  // Arithmetic operators
254  m_Pointer += (difference_type)1;
255  return *this;
256  }
258  multi_ptr result(*this);
259  ++(*this);
260  return result;
261  }
263  m_Pointer -= (difference_type)1;
264  return *this;
265  }
267  multi_ptr result(*this);
268  --(*this);
269  return result;
270  }
272  m_Pointer += r;
273  return *this;
274  }
276  m_Pointer -= r;
277  return *this;
278  }
280  return multi_ptr(m_Pointer + r);
281  }
283  return multi_ptr(m_Pointer - r);
284  }
285 
286 #ifdef __ENABLE_USM_ADDR_SPACE__
287  // Explicit conversion to global_space
288  // Only available if Space == address_space::global_device_space ||
289  // Space == address_space::global_host_space
290  template <access::address_space _Space = Space,
291  typename = typename detail::enable_if_t<
292  _Space == Space &&
293  (Space == access::address_space::global_device_space ||
294  Space == access::address_space::global_host_space)>>
295  explicit
297  using global_pointer_t = typename detail::DecoratedType<
298  ElementType, access::address_space::global_space>::type *;
300  reinterpret_cast<global_pointer_t>(m_Pointer));
301  }
302 #endif // __ENABLE_USM_ADDR_SPACE__
303 
304  // Only if Space == global_space
305  template <
306  access::address_space _Space = Space,
307  typename = typename detail::enable_if_t<
308  _Space == Space && Space == access::address_space::global_space>>
309  void prefetch(size_t NumElements) const {
310  size_t NumBytes = NumElements * sizeof(ElementType);
311  using ptr_t = typename detail::DecoratedType<char, Space>::type const *;
312  __spirv_ocl_prefetch(reinterpret_cast<ptr_t>(m_Pointer), NumBytes);
313  }
314 
315 private:
316  pointer_t m_Pointer;
317 };
318 
319 // Specialization of multi_ptr for void
320 template <access::address_space Space> class multi_ptr<void, Space> {
321 public:
322  using element_type = void;
323  using difference_type = std::ptrdiff_t;
324 
325  // Implementation defined pointer types that correspond to
326  // SYCL/OpenCL interoperability types for OpenCL C functions
328  using const_pointer_t =
330 
331  static constexpr access::address_space address_space = Space;
332 
333  // Constructors
334  multi_ptr() : m_Pointer(nullptr) {}
335  multi_ptr(const multi_ptr &) = default;
336  multi_ptr(multi_ptr &&) = default;
337  multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
338 #ifdef __SYCL_DEVICE_ONLY__
339  multi_ptr(void *pointer) : m_Pointer((pointer_t)pointer) {
340  // TODO An implementation should reject an argument if the deduced
341  // address space is not compatible with Space.
342  }
343 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
344  template <typename = typename detail::const_if_const_AS<Space, void>>
345  multi_ptr(const void *pointer) : m_Pointer((pointer_t)(pointer)) {}
346 #endif
347 #endif
348  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
349  ~multi_ptr() = default;
350 
351  // TODO: This constructor is the temporary solution for the existing problem
352  // with conversions from multi_ptr<ElementType, Space> to
353  // multi_ptr<void, Space>. Without it the compiler
354  // fails due to having 3 different same rank paths available.
355  template <typename ElementType>
356  multi_ptr(const multi_ptr<ElementType, Space> &ETP) : m_Pointer(ETP.get()) {}
357 
358  // Assignment operators
359  multi_ptr &operator=(const multi_ptr &) = default;
360  multi_ptr &operator=(multi_ptr &&) = default;
362  m_Pointer = pointer;
363  return *this;
364  }
365 #ifdef __SYCL_DEVICE_ONLY__
366  multi_ptr &operator=(void *pointer) {
367  // TODO An implementation should reject an argument if the deduced
368  // address space is not compatible with Space.
369  m_Pointer = (pointer_t)pointer;
370  return *this;
371  }
372 #endif
373  multi_ptr &operator=(std::nullptr_t) {
374  m_Pointer = nullptr;
375  return *this;
376  }
377 
378  // Only if Space == global_space || global_device_space
379  template <typename ElementType, int dimensions, access::mode Mode,
380  typename PropertyListT, access::address_space _Space = Space,
381  typename = typename detail::enable_if_t<
382  _Space == Space &&
383  (Space == access::address_space::global_space ||
384  Space == access::address_space::global_device_space)>>
386  accessor<ElementType, dimensions, Mode, access::target::global_buffer,
387  access::placeholder::false_t, PropertyListT>
388  Accessor)
389  : multi_ptr(Accessor.get_pointer()) {}
390 
391  // Only if Space == local_space
392  template <typename ElementType, int dimensions, access::mode Mode,
393  typename PropertyListT, access::address_space _Space = Space,
394  typename = typename detail::enable_if_t<
395  _Space == Space && Space == access::address_space::local_space>>
396  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
397  access::placeholder::false_t, PropertyListT>
398  Accessor)
399  : multi_ptr(Accessor.get_pointer()) {}
400 
401  // Only if Space == constant_space
402  template <
403  typename ElementType, int dimensions, access::mode Mode,
404  typename PropertyListT, access::address_space _Space = Space,
405  typename = typename detail::enable_if_t<
406  _Space == Space && Space == access::address_space::constant_space>>
408  accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
409  access::placeholder::false_t, PropertyListT>
410  Accessor)
411  : multi_ptr(Accessor.get_pointer()) {}
412 
414  // Returns the underlying OpenCL C pointer
415  pointer_t get() const { return m_Pointer; }
416 
417  // Implicit conversion to the underlying pointer type
418  operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); };
419 
420  // Explicit conversion to a multi_ptr<ElementType>
421  template <typename ElementType>
422  explicit operator multi_ptr<ElementType, Space>() const {
423  using elem_pointer_t =
426  static_cast<elem_pointer_t>(m_Pointer));
427  }
428 
429  // Implicit conversion to multi_ptr<const void, Space>
430  operator multi_ptr<const void, Space>() const {
431  using ptr_t = typename detail::DecoratedType<const void, Space>::type *;
432  return multi_ptr<const void, Space>(reinterpret_cast<ptr_t>(m_Pointer));
433  }
434 
435 private:
436  pointer_t m_Pointer;
437 };
438 
439 // Specialization of multi_ptr for const void
440 template <access::address_space Space>
441 class multi_ptr<const void, Space> {
442 public:
443  using element_type = const void;
444  using difference_type = std::ptrdiff_t;
445 
446  // Implementation defined pointer types that correspond to
447  // SYCL/OpenCL interoperability types for OpenCL C functions
449  using const_pointer_t =
451 
452  static constexpr access::address_space address_space = Space;
453 
454  // Constructors
455  multi_ptr() : m_Pointer(nullptr) {}
456  multi_ptr(const multi_ptr &) = default;
457  multi_ptr(multi_ptr &&) = default;
458  multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
459 #ifdef __SYCL_DEVICE_ONLY__
460  multi_ptr(const void *pointer) : m_Pointer((pointer_t)pointer) {
461  // TODO An implementation should reject an argument if the deduced
462  // address space is not compatible with Space.
463  }
464 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
465  template <typename = typename detail::const_if_const_AS<Space, void>>
466  multi_ptr(const void *pointer) : m_Pointer((pointer_t)(pointer)) {}
467 #endif
468 #endif
469  multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
470  ~multi_ptr() = default;
471 
472  // TODO: This constructor is the temporary solution for the existing problem
473  // with conversions from multi_ptr<ElementType, Space> to
474  // multi_ptr<const void, Space>. Without it the compiler
475  // fails due to having 3 different same rank paths available.
476  template <typename ElementType>
477  multi_ptr(const multi_ptr<ElementType, Space> &ETP) : m_Pointer(ETP.get()) {}
478 
479  // Assignment operators
480  multi_ptr &operator=(const multi_ptr &) = default;
481  multi_ptr &operator=(multi_ptr &&) = default;
483  m_Pointer = pointer;
484  return *this;
485  }
486 #ifdef __SYCL_DEVICE_ONLY__
487  multi_ptr &operator=(const void *pointer) {
488  // TODO An implementation should reject an argument if the deduced
489  // address space is not compatible with Space.
490  m_Pointer = (pointer_t)pointer;
491  return *this;
492  }
493 #endif
494  multi_ptr &operator=(std::nullptr_t) {
495  m_Pointer = nullptr;
496  return *this;
497  }
498 
499  // Only if Space == global_space || global_device_space
500  template <typename ElementType, int dimensions, access::mode Mode,
501  typename PropertyListT, access::address_space _Space = Space,
502  typename = typename detail::enable_if_t<
503  _Space == Space &&
504  (Space == access::address_space::global_space ||
505  Space == access::address_space::global_device_space)>>
507  accessor<ElementType, dimensions, Mode, access::target::global_buffer,
508  access::placeholder::false_t, PropertyListT>
509  Accessor)
510  : multi_ptr(Accessor.get_pointer()) {}
511 
512  // Only if Space == local_space
513  template <typename ElementType, int dimensions, access::mode Mode,
514  typename PropertyListT, access::address_space _Space = Space,
515  typename = typename detail::enable_if_t<
516  _Space == Space && Space == access::address_space::local_space>>
517  multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
518  access::placeholder::false_t, PropertyListT>
519  Accessor)
520  : multi_ptr(Accessor.get_pointer()) {}
521 
522  // Only if Space == constant_space
523  template <
524  typename ElementType, int dimensions, access::mode Mode,
525  typename PropertyListT, access::address_space _Space = Space,
526  typename = typename detail::enable_if_t<
527  _Space == Space && Space == access::address_space::constant_space>>
529  accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
530  access::placeholder::false_t, PropertyListT>
531  Accessor)
532  : multi_ptr(Accessor.get_pointer()) {}
533 
534  // Returns the underlying OpenCL C pointer
535  pointer_t get() const { return m_Pointer; }
536 
537  // Implicit conversion to the underlying pointer type
538  operator const void*() const {
539  return reinterpret_cast<const void *>(m_Pointer);
540  };
541 
542  // Explicit conversion to a multi_ptr<const ElementType>
543  // multi_ptr<const void, Space> -> multi_ptr<const void, Space>
544  // The result type must have const specifier.
545  template <typename ElementType>
546  explicit operator multi_ptr<const ElementType, Space>() const {
547  using elem_pointer_t =
550  static_cast<elem_pointer_t>(m_Pointer));
551  }
552 
553 private:
554  pointer_t m_Pointer;
555 };
556 
557 #ifdef __cpp_deduction_guides
558 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
559  typename PropertyListT, class T>
560 multi_ptr(accessor<T, dimensions, Mode, access::target::global_buffer,
561  isPlaceholder, PropertyListT>)
562  ->multi_ptr<T, access::address_space::global_space>;
563 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
564  typename PropertyListT, class T>
565 multi_ptr(accessor<T, dimensions, Mode, access::target::constant_buffer,
566  isPlaceholder, PropertyListT>)
567  ->multi_ptr<T, access::address_space::constant_space>;
568 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
569  typename PropertyListT, class T>
570 multi_ptr(accessor<T, dimensions, Mode, access::target::local, isPlaceholder,
571  PropertyListT>)
572  ->multi_ptr<T, access::address_space::local_space>;
573 #endif
574 
575 template <typename ElementType, access::address_space Space>
576 multi_ptr<ElementType, Space>
578  return multi_ptr<ElementType, Space>(pointer);
579 }
580 
581 #ifdef __SYCL_DEVICE_ONLY__
582 // An implementation should reject an argument if the deduced address space
583 // is not compatible with Space.
584 // This is guaranteed by the c'tor.
585 template <typename ElementType, access::address_space Space>
586 multi_ptr<ElementType, Space> make_ptr(ElementType *pointer) {
587  return multi_ptr<ElementType, Space>(pointer);
588 }
589 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
590 template <typename ElementType, access::address_space Space,
591  typename = typename detail::const_if_const_AS<Space, ElementType>>
592 multi_ptr<ElementType, Space> make_ptr(const ElementType *pointer) {
593  return multi_ptr<ElementType, Space>(pointer);
594 }
595 #endif // RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR
596 #endif // // __SYCL_DEVICE_ONLY__
597 
598 template <typename ElementType, access::address_space Space>
600  const multi_ptr<ElementType, Space> &rhs) {
601  return lhs.get() == rhs.get();
602 }
603 
604 template <typename ElementType, access::address_space Space>
606  const multi_ptr<ElementType, Space> &rhs) {
607  return lhs.get() != rhs.get();
608 }
609 
610 template <typename ElementType, access::address_space Space>
612  const multi_ptr<ElementType, Space> &rhs) {
613  return lhs.get() < rhs.get();
614 }
615 
616 template <typename ElementType, access::address_space Space>
618  const multi_ptr<ElementType, Space> &rhs) {
619  return lhs.get() > rhs.get();
620 }
621 
622 template <typename ElementType, access::address_space Space>
624  const multi_ptr<ElementType, Space> &rhs) {
625  return lhs.get() <= rhs.get();
626 }
627 
628 template <typename ElementType, access::address_space Space>
630  const multi_ptr<ElementType, Space> &rhs) {
631  return lhs.get() >= rhs.get();
632 }
633 
634 template <typename ElementType, access::address_space Space>
635 bool operator!=(const multi_ptr<ElementType, Space> &lhs, std::nullptr_t) {
636  return lhs.get() != nullptr;
637 }
638 
639 template <typename ElementType, access::address_space Space>
640 bool operator!=(std::nullptr_t, const multi_ptr<ElementType, Space> &rhs) {
641  return rhs.get() != nullptr;
642 }
643 
644 template <typename ElementType, access::address_space Space>
645 bool operator==(const multi_ptr<ElementType, Space> &lhs, std::nullptr_t) {
646  return lhs.get() == nullptr;
647 }
648 
649 template <typename ElementType, access::address_space Space>
650 bool operator==(std::nullptr_t, const multi_ptr<ElementType, Space> &rhs) {
651  return rhs.get() == nullptr;
652 }
653 
654 template <typename ElementType, access::address_space Space>
655 bool operator>(const multi_ptr<ElementType, Space> &lhs, std::nullptr_t) {
656  return lhs.get() != nullptr;
657 }
658 
659 template <typename ElementType, access::address_space Space>
660 bool operator>(std::nullptr_t, const multi_ptr<ElementType, Space> &) {
661  return false;
662 }
663 
664 template <typename ElementType, access::address_space Space>
665 bool operator<(const multi_ptr<ElementType, Space> &, std::nullptr_t) {
666  return false;
667 }
668 
669 template <typename ElementType, access::address_space Space>
670 bool operator<(std::nullptr_t, const multi_ptr<ElementType, Space> &rhs) {
671  return rhs.get() != nullptr;
672 }
673 
674 template <typename ElementType, access::address_space Space>
675 bool operator>=(const multi_ptr<ElementType, Space> &, std::nullptr_t) {
676  return true;
677 }
678 
679 template <typename ElementType, access::address_space Space>
680 bool operator>=(std::nullptr_t, const multi_ptr<ElementType, Space> &rhs) {
681  return rhs.get() == nullptr;
682 }
683 
684 template <typename ElementType, access::address_space Space>
685 bool operator<=(const multi_ptr<ElementType, Space> &lhs, std::nullptr_t) {
686  return lhs.get() == nullptr;
687 }
688 
689 template <typename ElementType, access::address_space Space>
690 bool operator<=(std::nullptr_t, const multi_ptr<ElementType, Space> &rhs) {
691  return rhs.get() == nullptr;
692 }
693 
694 } // namespace sycl
695 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::multi_ptr< void, Space >::multi_ptr
multi_ptr()
Definition: multi_ptr.hpp:334
spirv_ops.hpp
cl::sycl::multi_ptr::const_pointer_t
typename detail::DecoratedType< ElementType, Space >::type const * const_pointer_t
Definition: multi_ptr.hpp:41
cl::sycl::multi_ptr::multi_ptr
multi_ptr()
Definition: multi_ptr.hpp:50
cl::sycl::multi_ptr::operator+=
multi_ptr & operator+=(difference_type r)
Definition: multi_ptr.hpp:271
cl::sycl::multi_ptr< void, Space >::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::constant_buffer, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:407
cl::sycl::multi_ptr< const void, Space >::operator=
multi_ptr & operator=(std::nullptr_t)
Definition: multi_ptr.hpp:494
type
cl::sycl::multi_ptr::operator[]
ReturnRef operator[](difference_type index)
Definition: multi_ptr.hpp:104
cl::sycl::multi_ptr< void, Space >::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::global_buffer, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:385
cl::sycl::multi_ptr::multi_ptr
multi_ptr(ElementType *pointer)
Definition: multi_ptr.hpp:57
type_traits.hpp
cl::sycl::multi_ptr::operator[]
ReturnConstRef operator[](difference_type index) const
Definition: multi_ptr.hpp:108
cl::sycl::access::placeholder
placeholder
Definition: access.hpp:43
cl::sycl::multi_ptr::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:133
cl::sycl::multi_ptr::get
pointer_t get() const
Definition: multi_ptr.hpp:215
cl::sycl::multi_ptr::multi_ptr
multi_ptr(accessor< typename detail::remove_const_t< ET >, dimensions, Mode, access::target::global_buffer, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:169
cl::sycl::multi_ptr< const void, Space >::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::local, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:517
cl::sycl::multi_ptr::pointer_t
typename detail::DecoratedType< ElementType, Space >::type * pointer_t
Definition: multi_ptr.hpp:39
cl::sycl::multi_ptr< const void, Space >::operator=
multi_ptr & operator=(pointer_t pointer)
Definition: multi_ptr.hpp:482
cl::sycl::operator<=
bool operator<=(std::nullptr_t, const multi_ptr< ElementType, Space > &rhs)
Definition: multi_ptr.hpp:690
cl::sycl::detail::half_impl::BIsRepresentationT
half BIsRepresentationT
Definition: half_type.hpp:301
cl::sycl::multi_ptr::ReturnPtr
detail::const_if_const_AS< Space, ElementType > * ReturnPtr
Definition: multi_ptr.hpp:92
cl::sycl::multi_ptr< const void, Space >::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::constant_buffer, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:528
__spirv_ocl_prefetch
void __spirv_ocl_prefetch(const char *Ptr, size_t NumBytes) noexcept
Definition: spirv_ops.cpp:47
cl::sycl::multi_ptr::operator--
multi_ptr operator--(int)
Definition: multi_ptr.hpp:266
cl::sycl::multi_ptr< const void, Space >::multi_ptr
multi_ptr()
Definition: multi_ptr.hpp:455
cl::sycl::multi_ptr::operator++
multi_ptr operator++(int)
Definition: multi_ptr.hpp:257
cl::sycl::multi_ptr< void, Space >::multi_ptr
multi_ptr(std::nullptr_t)
Definition: multi_ptr.hpp:348
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:32
cl::sycl::multi_ptr::const_reference_t
typename detail::DecoratedType< ElementType, Space >::type & const_reference_t
Definition: multi_ptr.hpp:45
cl::sycl::multi_ptr< const void, Space >::const_pointer_t
typename detail::DecoratedType< const void, Space >::type const * const_pointer_t
Definition: multi_ptr.hpp:450
access.hpp
operator==
bool operator==(const Slab &Lhs, const Slab &Rhs)
Definition: usm_allocator.cpp:343
cl::sycl::multi_ptr< const void, Space >::element_type
const void element_type
Definition: multi_ptr.hpp:443
cl::sycl::detail::const_if_const_AS
DataT const_if_const_AS
Definition: type_traits.hpp:351
cl::sycl::multi_ptr::prefetch
void prefetch(size_t NumElements) const
Definition: multi_ptr.hpp:309
cl::sycl::operator!=
bool operator!=(std::nullptr_t, const multi_ptr< ElementType, Space > &rhs)
Definition: multi_ptr.hpp:640
cl::sycl::detail::get
Definition: tuple.hpp:59
cl::sycl::multi_ptr::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 >> &ETP)
Definition: multi_ptr.hpp:209
cl::sycl::multi_ptr< void, Space >
Definition: multi_ptr.hpp:320
cl::sycl::operator>
bool operator>(std::nullptr_t, const multi_ptr< ElementType, Space > &)
Definition: multi_ptr.hpp:660
cl::sycl::multi_ptr::multi_ptr
multi_ptr(accessor< typename detail::remove_const_t< ET >, dimensions, Mode, access::target::local, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:183
cl::sycl::multi_ptr::ReturnConstRef
const ElementType & ReturnConstRef
Definition: multi_ptr.hpp:94
cl::sycl::multi_ptr< void, Space >::pointer_t
typename detail::DecoratedType< void, Space >::type * pointer_t
Definition: multi_ptr.hpp:327
cl::sycl::multi_ptr< void, Space >::ReturnPtr
detail::const_if_const_AS< Space, void > * ReturnPtr
Definition: multi_ptr.hpp:413
cl::sycl::multi_ptr< void, Space >::operator=
multi_ptr & operator=(pointer_t pointer)
Definition: multi_ptr.hpp:361
cl::sycl::multi_ptr< const void, Space >::get
pointer_t get() const
Definition: multi_ptr.hpp:535
cl::sycl::multi_ptr< const void, Space >::difference_type
std::ptrdiff_t difference_type
Definition: multi_ptr.hpp:444
cl::sycl::operator>=
bool operator>=(std::nullptr_t, const multi_ptr< ElementType, Space > &rhs)
Definition: multi_ptr.hpp:680
cl::sycl::multi_ptr::operator=
multi_ptr & operator=(ElementType *pointer)
Definition: multi_ptr.hpp:80
cl::sycl::multi_ptr::operator++
multi_ptr & operator++()
Definition: multi_ptr.hpp:253
cl::sycl::multi_ptr::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::constant_buffer, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:144
cl::sycl::accessor
Buffer accessor.
Definition: accessor.hpp:225
cl::sycl::multi_ptr< void, Space >::operator=
multi_ptr & operator=(std::nullptr_t)
Definition: multi_ptr.hpp:373
cl::sycl::access::target
target
Definition: access.hpp:17
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::multi_ptr::reference_t
typename detail::DecoratedType< ElementType, Space >::type & reference_t
Definition: multi_ptr.hpp:43
cl::sycl::multi_ptr::operator-=
multi_ptr & operator-=(difference_type r)
Definition: multi_ptr.hpp:275
cl::sycl::image_channel_order::r
@ r
cl::sycl::multi_ptr< void, Space >::get
pointer_t get() const
Definition: multi_ptr.hpp:415
cl::sycl::multi_ptr< const void, Space >
Definition: multi_ptr.hpp:441
cl::sycl::multi_ptr::ReturnRef
detail::const_if_const_AS< Space, ElementType > & ReturnRef
Definition: multi_ptr.hpp:93
cl::sycl::make_ptr
multi_ptr< ElementType, Space > make_ptr(typename multi_ptr< ElementType, Space >::pointer_t pointer)
Definition: multi_ptr.hpp:577
cl::sycl::multi_ptr::operator=
multi_ptr & operator=(std::nullptr_t)
Definition: multi_ptr.hpp:87
cl::sycl::access::address_space
address_space
Definition: access.hpp:45
cl::sycl::multi_ptr< void, Space >::multi_ptr
multi_ptr(const multi_ptr< ElementType, Space > &ETP)
Definition: multi_ptr.hpp:356
cl::sycl::multi_ptr< void, Space >::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::local, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:396
cl::sycl::multi_ptr::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::global_buffer, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:120
cl::sycl::multi_ptr::operator-
multi_ptr operator-(difference_type r) const
Definition: multi_ptr.hpp:282
cl::sycl::multi_ptr< const void, Space >::multi_ptr
multi_ptr(accessor< ElementType, dimensions, Mode, access::target::global_buffer, access::placeholder::false_t, PropertyListT > Accessor)
Definition: multi_ptr.hpp:506
cl::sycl::multi_ptr< void, Space >::const_pointer_t
typename detail::DecoratedType< void, Space >::type const * const_pointer_t
Definition: multi_ptr.hpp:329
cl::sycl::multi_ptr< const void, Space >::pointer_t
typename detail::DecoratedType< const void, Space >::type * pointer_t
Definition: multi_ptr.hpp:448
cl::sycl::multi_ptr< void, Space >::multi_ptr
multi_ptr(pointer_t pointer)
Definition: multi_ptr.hpp:337
cl::sycl::multi_ptr< const void, Space >::multi_ptr
multi_ptr(const multi_ptr< ElementType, Space > &ETP)
Definition: multi_ptr.hpp:477
cl::sycl::multi_ptr< const void, Space >::multi_ptr
multi_ptr(pointer_t pointer)
Definition: multi_ptr.hpp:458
cl::sycl::detail::remove_const_t
typename std::remove_const< T >::type remove_const_t
Definition: stl_type_traits.hpp:30
cl::sycl::multi_ptr::operator+
multi_ptr operator+(difference_type r) const
Definition: multi_ptr.hpp:279
common.hpp
cl::sycl::multi_ptr::multi_ptr
multi_ptr(accessor< typename detail::remove_const_t< ET >, dimensions, Mode, access::target::constant_buffer, isPlaceholder, PropertyListT > Accessor)
Definition: multi_ptr.hpp:196
cl::sycl::access::mode
mode
Definition: access.hpp:28
cl::sycl::multi_ptr::operator*
ReturnRef operator*() const
Definition: multi_ptr.hpp:96
cl::sycl::multi_ptr< void, Space >::difference_type
std::ptrdiff_t difference_type
Definition: multi_ptr.hpp:323
cl::sycl::multi_ptr< const void, Space >::multi_ptr
multi_ptr(std::nullptr_t)
Definition: multi_ptr.hpp:469
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::operator<
bool operator<(std::nullptr_t, const multi_ptr< ElementType, Space > &rhs)
Definition: multi_ptr.hpp:670
cl::sycl::multi_ptr::difference_type
std::ptrdiff_t difference_type
Definition: multi_ptr.hpp:35
cl::sycl::multi_ptr::operator--
multi_ptr & operator--()
Definition: multi_ptr.hpp:262
cl::sycl::multi_ptr< void, Space >::element_type
void element_type
Definition: multi_ptr.hpp:322
cl::sycl::multi_ptr::operator->
ReturnPtr operator->() const
Definition: multi_ptr.hpp:100
cl::sycl::multi_ptr::multi_ptr
multi_ptr(std::nullptr_t)
Definition: multi_ptr.hpp:66
cl::sycl::errc::accessor
@ accessor
cl::sycl::detail::DecoratedType
Definition: access.hpp:157
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12