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