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