DPC++ Runtime
Runtime libraries for oneAPI DPC++
annotated_ptr.hpp
Go to the documentation of this file.
1 //
2 //==----------- annotated_ptr.hpp - SYCL annotated_ptr extension -----------==//
3 //
4 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5 // See https://llvm.org/LICENSE.txt for license information.
6 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7 //
8 //===----------------------------------------------------------------------===//
9 
10 #pragma once
11 
12 #include <sycl/detail/defines.hpp>
20 
21 #include <cstddef>
22 #include <string_view>
23 #include <tuple>
24 #include <type_traits>
25 #include <utility>
26 #include <variant>
27 
28 namespace sycl {
29 inline namespace _V1 {
30 namespace ext {
31 namespace oneapi {
32 namespace experimental {
33 
34 template <typename T, typename PropertyListT = empty_properties_t>
36  // This should always fail when instantiating the unspecialized version.
37  static constexpr bool is_valid_property_list =
39  static_assert(is_valid_property_list, "Property list is invalid.");
40 };
41 
42 namespace detail {
43 template <class T> struct is_ann_ref_impl : std::false_type {};
44 template <class T, class P>
45 struct is_ann_ref_impl<annotated_ref<T, P>> : std::true_type {};
46 template <class T, class P>
47 struct is_ann_ref_impl<const annotated_ref<T, P>> : std::true_type {};
48 template <class T>
49 constexpr bool is_ann_ref_v =
51 
52 template <typename... Ts>
54  detail::ContainsProperty<alignment_key, std::tuple<Ts...>>;
55 
56 // properties filter
57 template <typename property_list, template <class...> typename filter>
59  sycl::detail::boost::mp11::mp_copy_if<property_list, filter>;
60 
61 // filter properties that are applied on annotations
62 template <typename... Props>
64  PropertiesFilter<std::tuple<Props...>, propagateToPtrAnnotation>>;
65 } // namespace detail
66 
67 template <typename I, typename P> struct annotationHelper {};
68 
69 // unpack properties to varadic template
70 template <typename I, typename... P>
71 struct annotationHelper<I, detail::properties_t<P...>> {
72  static I *annotate(I *ptr) {
73  return __builtin_intel_sycl_ptr_annotation(
76  }
77 
78  static I load(I *ptr) {
79  return *__builtin_intel_sycl_ptr_annotation(
82  }
83 
84  template <class O> static I store(I *ptr, O &&Obj) {
85  return *__builtin_intel_sycl_ptr_annotation(
87  detail::PropertyMetaInfo<P>::value...) = std::forward<O>(Obj);
88  }
89 };
90 
91 template <typename T, typename... Props>
92 class annotated_ref<T, detail::properties_t<Props...>> {
93  using property_list_t = detail::properties_t<Props...>;
94 
95  static_assert(
96  std::is_trivially_copyable_v<T>,
97  "annotated_ref can only encapsulate a trivially-copyable type!");
98 
99 private:
100  T *m_Ptr;
101  explicit annotated_ref(T *Ptr) : m_Ptr(Ptr) {}
102 
103 public:
104  annotated_ref(const annotated_ref &) = delete;
105 
106  // implicit conversion with annotaion
107  operator T() const {
108 #ifdef __SYCL_DEVICE_ONLY__
109  return annotationHelper<T, detail::annotation_filter<Props...>>::load(
110  m_Ptr);
111 #else
112  return *m_Ptr;
113 #endif
114  }
115 
116  // assignment operator with annotaion
117  template <class O, typename = std::enable_if_t<!detail::is_ann_ref_v<O>>>
118  T operator=(O &&Obj) const {
119 #ifdef __SYCL_DEVICE_ONLY__
120  return annotationHelper<T, detail::annotation_filter<Props...>>::store(
121  m_Ptr, Obj);
122 #else
123  return *m_Ptr = std::forward<O>(Obj);
124 #endif
125  }
126 
127  template <class O, class P>
128  T operator=(const annotated_ref<O, P> &Ref) const {
129  O t2 = Ref.operator O();
130  return *this = t2;
131  }
132 
133  // propagate compound operators
134 #define PROPAGATE_OP(op) \
135  template <class O, typename = std::enable_if_t<!detail::is_ann_ref_v<O>>> \
136  T operator op(O &&rhs) const { \
137  T t = this->operator T(); \
138  t op std::forward<O>(rhs); \
139  *this = t; \
140  return t; \
141  } \
142  template <class O, class P> \
143  T operator op(const annotated_ref<O, P> &rhs) const { \
144  T t = this->operator T(); \
145  O t2 = rhs.operator T(); \
146  t op t2; \
147  *this = t; \
148  return t; \
149  }
150  PROPAGATE_OP(+=)
151  PROPAGATE_OP(-=)
152  PROPAGATE_OP(*=)
153  PROPAGATE_OP(/=)
154  PROPAGATE_OP(%=)
155  PROPAGATE_OP(^=)
156  PROPAGATE_OP(&=)
157  PROPAGATE_OP(|=)
158  PROPAGATE_OP(<<=)
159  PROPAGATE_OP(>>=)
160 #undef PROPAGATE_OP
161 
162  // propagate binary operators
163 #define PROPAGATE_OP(op) \
164  template <class O> \
165  friend auto operator op(O &&a, const annotated_ref &b) \
166  ->decltype(std::forward<O>(a) op std::declval<T>()) { \
167  return std::forward<O>(a) op b.operator T(); \
168  } \
169  template <class O, typename = std::enable_if_t<!detail::is_ann_ref_v<O>>> \
170  friend auto operator op(const annotated_ref &a, O &&b) \
171  ->decltype(std::declval<T>() op std::forward<O>(b)) { \
172  return a.operator T() op std::forward<O>(b); \
173  }
174  PROPAGATE_OP(+)
175  PROPAGATE_OP(-)
176  PROPAGATE_OP(*)
177  PROPAGATE_OP(/)
178  PROPAGATE_OP(%)
179  PROPAGATE_OP(|)
180  PROPAGATE_OP(&)
181  PROPAGATE_OP(^)
182  PROPAGATE_OP(<<)
183  PROPAGATE_OP(>>)
184  PROPAGATE_OP(<)
185  PROPAGATE_OP(<=)
186  PROPAGATE_OP(>)
187  PROPAGATE_OP(>=)
188  PROPAGATE_OP(==)
189  PROPAGATE_OP(!=)
190  PROPAGATE_OP(&&)
191  PROPAGATE_OP(||)
192 #undef PROPAGATE_OP
193 
194 // Propagate unary operators
195 // by setting a default template we get SFINAE to kick in
196 #define PROPAGATE_OP(op) \
197  template <typename O = T> \
198  auto operator op() const->decltype(op std::declval<O>()) { \
199  return op this->operator O(); \
200  }
201  PROPAGATE_OP(+)
202  PROPAGATE_OP(-)
203  PROPAGATE_OP(!)
204  PROPAGATE_OP(~)
205 #undef PROPAGATE_OP
206 
207  // Propagate inc/dec operators
208  T operator++() const {
209  T t = this->operator T();
210  ++t;
211  *this = t;
212  return t;
213  }
214 
215  T operator++(int) const {
216  T t1 = this->operator T();
217  T t2 = t1;
218  t2++;
219  *this = t2;
220  return t1;
221  }
222 
223  T operator--() const {
224  T t = this->operator T();
225  --t;
226  *this = t;
227  return t;
228  }
229 
230  T operator--(int) const {
231  T t1 = this->operator T();
232  T t2 = t1;
233  t2--;
234  *this = t2;
235  return t1;
236  }
237 
238  template <class T2, class P2> friend class annotated_ptr;
239 };
240 
241 #ifdef __cpp_deduction_guides
242 template <typename T, typename... Args>
243 annotated_ptr(T *, Args...)
244  -> annotated_ptr<T, typename detail::DeducedProperties<Args...>::type>;
245 
246 template <typename T, typename old, typename... ArgT>
247 annotated_ptr(annotated_ptr<T, old>, properties<std::tuple<ArgT...>>)
248  -> annotated_ptr<
250 #endif // __cpp_deduction_guides
251 
252 template <typename T, typename PropertyListT = empty_properties_t>
254  // This should always fail when instantiating the unspecialized version.
255  static constexpr bool is_valid_property_list =
257  static_assert(is_valid_property_list, "Property list is invalid.");
258 };
259 
260 template <typename T, typename... Props>
262 __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
263 
264  static_assert(std::is_same_v<T, void> || std::is_trivially_copyable_v<T>,
265  "annotated_ptr can only encapsulate either "
266  "a trivially-copyable type "
267  "or void!");
268 
270 
271  // annotated_ref type
272  using reference =
274 
275 #ifdef __ENABLE_USM_ADDR_SPACE__
276  using global_pointer_t = std::conditional_t<
279  std::conditional_t<
283 #else
285 #endif // __ENABLE_USM_ADDR_SPACE__
286 
287  T *m_Ptr;
288 
289  template <typename T2, typename PropertyListT> friend class annotated_ptr;
290 
291 #ifdef __SYCL_DEVICE_ONLY__
292  void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter(
295  m_Ptr = Ptr;
296  }
297 #endif
298 
299 public:
300  annotated_ptr() noexcept = default;
301  annotated_ptr(const annotated_ptr &) = default;
302  annotated_ptr &operator=(const annotated_ptr &) = default;
303 
304  explicit annotated_ptr(T *Ptr,
305  const property_list_t & = properties{}) noexcept
306  : m_Ptr(Ptr) {}
307 
308  // Constructs an annotated_ptr object from a raw pointer and variadic
309  // properties. The new property set contains all properties of the input
310  // variadic properties. The same property in `Props...` and
311  // `PropertyValueTs...` must have the same property value.
312  template <typename... PropertyValueTs>
313  explicit annotated_ptr(T *Ptr, const PropertyValueTs &...props) noexcept
314  : m_Ptr(Ptr) {
315 
316  static constexpr bool has_same_properties = std::is_same<
319  decltype(properties{props...})>>::value;
320  static_assert(
321  has_same_properties,
322  "The property list must contain all properties of the input of the "
323  "constructor");
324  }
325 
326  // Constructs an annotated_ptr object from another annotated_ptr object.
327  // The new property set contains all properties of the input
328  // annotated_ptr object. The same property in `Props...` and `PropertyList2`
329  // must have the same property value.
330  template <typename T2, typename PropertyList2>
332  : m_Ptr(other.m_Ptr) {
333  static constexpr bool is_input_convertible =
334  std::is_convertible<T2 *, T *>::value;
335  static_assert(
336  is_input_convertible,
337  "The underlying pointer type of the input annotated_ptr is not "
338  "convertible to the target pointer type");
339 
340  static constexpr bool has_same_properties = std::is_same<
343  static_assert(
344  has_same_properties,
345  "The constructed annotated_ptr type must contain all the properties "
346  "of the input annotated_ptr");
347  }
348 
349  // Constructs an annotated_ptr object from another annotated_ptr object and
350  // a property list. The new property set is the union of property lists
351  // `PropertyListU` and `PropertyListV`. The same property in `PropertyListU`
352  // and `PropertyListV` must have the same property value.
353  template <typename T2, typename PropertyListU, typename PropertyListV>
355  const PropertyListV &) noexcept
356  : m_Ptr(other.m_Ptr) {
357  static constexpr bool is_input_convertible =
358  std::is_convertible<T2 *, T *>::value;
359  static_assert(
360  is_input_convertible,
361  "The underlying pointer type of the input annotated_ptr is not "
362  "convertible to the target pointer type");
363 
364  static constexpr bool has_same_properties = std::is_same<
367  static_assert(
368  has_same_properties,
369  "The property list of constructed annotated_ptr type must be the "
370  "union of the input property lists");
371  }
372 
374 
375  std::ptrdiff_t operator-(annotated_ptr other) const noexcept {
376  return m_Ptr - other.m_Ptr;
377  }
378 
379  explicit operator bool() const noexcept { return m_Ptr != nullptr; }
380 
381  operator T *() const noexcept = delete;
382 
383  T *get() const noexcept {
384 #ifdef __SYCL_DEVICE_ONLY__
385  return annotationHelper<T, detail::annotation_filter<Props...>>::annotate(
386  m_Ptr);
387 #else
388  return m_Ptr;
389 #endif
390  }
391 
392  // When the properties contain alignment, operator '[]', '+', '++' and '--'
393  // (both post- and prefix) are disabled. Calling these operators when
394  // alignment is present causes a compile error. Note that clang format is
395  // turned off for these operators to make sure the complete error notes are
396  // printed
397  // clang-format off
398  template <bool has_alignment = detail::contains_alignment<Props...>::value,
399  class = std::enable_if_t<!has_alignment>>
400  reference operator[](std::ptrdiff_t idx) const noexcept {
401  return reference(m_Ptr + idx);
402  }
403 
404  template <bool has_alignment = detail::contains_alignment<Props...>::value,
405  class = std::enable_if_t<has_alignment>>
406  auto operator[](std::ptrdiff_t idx) const noexcept -> decltype("operator[] is not available when alignment is specified!") = delete;
407 
408  template <bool has_alignment = detail::contains_alignment<Props...>::value,
409  class = std::enable_if_t<!has_alignment>>
410  annotated_ptr operator+(size_t offset) const noexcept {
411  return annotated_ptr(m_Ptr + offset);
412  }
413 
414  template <bool has_alignment = detail::contains_alignment<Props...>::value,
415  class = std::enable_if_t<has_alignment>>
416  auto operator+(size_t offset) const noexcept -> decltype("operator+ is not available when alignment is specified!") = delete;
417 
418  template <bool has_alignment = detail::contains_alignment<Props...>::value,
419  class = std::enable_if_t<!has_alignment>>
421  m_Ptr += 1;
422  return *this;
423  }
424 
425  template <bool has_alignment = detail::contains_alignment<Props...>::value,
426  class = std::enable_if_t<has_alignment>>
427  auto operator++() noexcept -> decltype("operator++ is not available when alignment is specified!") = delete;
428 
429  template <bool has_alignment = detail::contains_alignment<Props...>::value,
430  class = std::enable_if_t<!has_alignment>>
431  annotated_ptr operator++(int) noexcept {
432  auto tmp = *this;
433  m_Ptr += 1;
434  return tmp;
435  }
436 
437  template <bool has_alignment = detail::contains_alignment<Props...>::value,
438  class = std::enable_if_t<has_alignment>>
439  auto operator++(int) noexcept -> decltype("operator++ is not available when alignment is specified!") = delete;
440 
441  template <bool has_alignment = detail::contains_alignment<Props...>::value,
442  class = std::enable_if_t<!has_alignment>>
444  m_Ptr -= 1;
445  return *this;
446  }
447 
448  template <bool has_alignment = detail::contains_alignment<Props...>::value,
449  class = std::enable_if_t<has_alignment>>
450  auto operator--() noexcept -> decltype("operator-- is not available when alignment is specified!") = delete;
451 
452  template <bool has_alignment = detail::contains_alignment<Props...>::value,
453  class = std::enable_if_t<!has_alignment>>
454  annotated_ptr operator--(int) noexcept {
455  auto tmp = *this;
456  m_Ptr -= 1;
457  return tmp;
458  }
459 
460  template <bool has_alignment = detail::contains_alignment<Props...>::value,
461  class = std::enable_if_t<has_alignment>>
462  auto operator--(int) noexcept -> decltype("operator-- is not available when alignment is specified!") = delete;
463 
464  // clang-format on
465 
466  template <typename PropertyT> static constexpr bool has_property() {
467  return property_list_t::template has_property<PropertyT>();
468  }
469 
470  template <typename PropertyT> static constexpr auto get_property() {
471  return property_list_t::template get_property<PropertyT>();
472  }
473 
474  // *************************************************************************
475  // All static error checking is added here instead of placing inside neat
476  // functions to minimize the number lines printed out when an assert
477  // is triggered.
478  // static constexprs are used to ensure that the triggered assert prints
479  // a message that is very readable. Without these, the assert will
480  // print out long templated names
481  // *************************************************************************
482  static constexpr bool is_valid_property_list =
484  static_assert(is_valid_property_list, "Property list is invalid.");
485  static constexpr bool contains_valid_properties =
486  check_property_list<T *, Props...>::value;
487  static_assert(contains_valid_properties,
488  "The property list contains invalid property.");
489  // check the set if FPGA specificed properties are used
490  static constexpr bool hasValidFPGAProperties =
491  detail::checkValidFPGAPropertySet<Props...>::value;
492  static_assert(hasValidFPGAProperties,
493  "FPGA Interface properties (i.e. awidth, dwidth, etc.) "
494  "can only be set with BufferLocation together.");
495  // check if conduit and register_map properties are specified together
496  static constexpr bool hasConduitAndRegisterMapProperties =
497  detail::checkHasConduitAndRegisterMap<Props...>::value;
499  "The properties conduit and register_map cannot be "
500  "specified at the same time.");
501 };
502 
503 } // namespace experimental
504 } // namespace oneapi
505 } // namespace ext
506 } // namespace _V1
507 } // namespace sycl
#define PROPAGATE_OP(op)
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
Definition: multi_ptr.hpp:95
Objects of the property_list class are containers for the SYCL properties.
#define __SYCL_SPECIAL_CLASS
Definition: defines.hpp:29
sycl::ext::oneapi::experimental::alignment_key alignment_key
The 'alignment' property is used to specify the alignment of memory accessed in ESIMD memory operatio...
struct sycl::detail::device_filter filter
sycl::detail::boost::mp11::mp_copy_if< property_list, filter > PropertiesFilter
typename merged_properties< LHSPropertiesT, RHSPropertiesT >::type merged_properties_t
Definition: properties.hpp:225
properties< std::tuple< PropertyValueTs... > > properties_t
Definition: properties.hpp:212
detail::ContainsProperty< alignment_key, std::tuple< Ts... > > contains_alignment
constexpr alignment_key::value_t< K > alignment
Definition: properties.hpp:74
annotated_ptr & operator++() noexcept
static constexpr bool has_property()
sycl::ext::oneapi::experimental::annotated_ref< T, property_list_t > reference
static constexpr bool hasValidFPGAProperties
std::ptrdiff_t operator-(annotated_ptr other) const noexcept
class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg< T *
reference operator*() const noexcept
static constexpr bool is_valid_property_list
static constexpr auto get_property()
detail::properties_t< Props... > property_list_t
static constexpr bool is_valid_property_list
static constexpr bool contains_valid_properties
annotated_ptr & operator--() noexcept
T & operator[](std::ptrdiff_t idx) const noexcept
annotated_ptr operator+(size_t offset) const noexcept
static constexpr bool hasConduitAndRegisterMapProperties
typename decorated_global_ptr< T >::pointer global_pointer_t
Definition: access.hpp:18
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324