@@ -53,8 +53,49 @@ template <typename... Ts>
53
53
using contains_alignment =
54
54
detail::ContainsProperty<alignment_key, std::tuple<Ts...>>;
55
55
56
+ // properties filter
57
+ template <typename property_list, template <class ...> typename filter>
58
+ using PropertiesFilter =
59
+ sycl::detail::boost::mp11::mp_copy_if<property_list, filter>;
60
+
61
+ // filter properties that are applied on annotations
62
+ template <typename ... Props>
63
+ using annotation_filter = properties<
64
+ PropertiesFilter<std::tuple<Props...>, propagateToPtrAnnotation>>;
56
65
} // namespace detail
57
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 (
74
+ ptr, detail::PropertyMetaInfo<P>::name...,
75
+ detail::PropertyMetaInfo<P>::value...);
76
+ }
77
+
78
+ // static I load(I *ptr) {
79
+ // return *annotate(ptr);
80
+ // }
81
+
82
+ // template <class O> static I store(I *ptr, O &&Obj) {
83
+ // return *annotate(ptr) = std::forward<O>(Obj);
84
+ // }
85
+
86
+ static I load (I *ptr) {
87
+ return *__builtin_intel_sycl_ptr_annotation (
88
+ ptr, detail::PropertyMetaInfo<P>::name...,
89
+ detail::PropertyMetaInfo<P>::value...);
90
+ }
91
+
92
+ template <class O > static I store (I *ptr, O &&Obj) {
93
+ return *__builtin_intel_sycl_ptr_annotation (
94
+ ptr, detail::PropertyMetaInfo<P>::name...,
95
+ detail::PropertyMetaInfo<P>::value...) = std::forward<O>(Obj);
96
+ }
97
+ };
98
+
58
99
template <typename T, typename ... Props>
59
100
class annotated_ref <T, detail::properties_t <Props...>> {
60
101
using property_list_t = detail::properties_t <Props...>;
@@ -67,44 +108,14 @@ class annotated_ref<T, detail::properties_t<Props...>> {
67
108
T *m_Ptr;
68
109
explicit annotated_ref (T *Ptr) : m_Ptr(Ptr) {}
69
110
70
- // properties filter
71
- template <typename property_list, template <class ...> typename filter>
72
- using PropertiesFilter =
73
- sycl::detail::boost::mp11::mp_copy_if<property_list, filter>;
74
-
75
- template <typename p>
76
- using annotation_filter = propagateToPtrAnnotation<typename p::key_t >;
77
-
78
- // filter properties that are applied on annotations
79
- using property_tuple_t = std::tuple<Props...>;
80
- using annotation_props =
81
- properties<PropertiesFilter<property_tuple_t , annotation_filter>>;
82
-
83
- template <typename I, typename P> struct annotationHelper {};
84
-
85
- // unpack properties to varadic template
86
- template <typename I, typename ... P>
87
- struct annotationHelper <I, detail::properties_t <P...>> {
88
- static I load (I *ptr) {
89
- return *__builtin_intel_sycl_ptr_annotation (
90
- ptr, detail::PropertyMetaInfo<P>::name...,
91
- detail::PropertyMetaInfo<P>::value...);
92
- }
93
-
94
- template <class O > static I store (I *ptr, O &&Obj) {
95
- return *__builtin_intel_sycl_ptr_annotation (
96
- ptr, detail::PropertyMetaInfo<P>::name...,
97
- detail::PropertyMetaInfo<P>::value...) = std::forward<O>(Obj);
98
- }
99
- };
100
-
101
111
public:
102
112
annotated_ref (const annotated_ref &) = delete ;
103
113
104
114
// implicit conversion with annotaion
105
115
operator T () const {
106
116
#ifdef __SYCL_DEVICE_ONLY__
107
- return annotationHelper<T, annotation_props>::load (m_Ptr);
117
+ return annotationHelper<T, detail::annotation_filter<Props...>>::load (
118
+ m_Ptr);
108
119
#else
109
120
return *m_Ptr;
110
121
#endif
@@ -114,7 +125,8 @@ class annotated_ref<T, detail::properties_t<Props...>> {
114
125
template <class O , typename = std::enable_if_t <!detail::is_ann_ref_v<O>>>
115
126
T operator =(O &&Obj) const {
116
127
#ifdef __SYCL_DEVICE_ONLY__
117
- return annotationHelper<T, annotation_props>::store (m_Ptr, Obj);
128
+ return annotationHelper<T, detail::annotation_filter<Props...>>::store (
129
+ m_Ptr, Obj);
118
130
#else
119
131
return *m_Ptr = std::forward<O>(Obj);
120
132
#endif
@@ -126,6 +138,16 @@ class annotated_ref<T, detail::properties_t<Props...>> {
126
138
return *this = t2;
127
139
}
128
140
141
+ // address-of operator
142
+ T *operator &() const {
143
+ #ifdef __SYCL_DEVICE_ONLY__
144
+ return annotationHelper<T, detail::annotation_filter<Props...>>::annotate (
145
+ m_Ptr);
146
+ #else
147
+ return *m_Ptr;
148
+ #endif
149
+ }
150
+
129
151
// propagate compound operators
130
152
#define PROPAGATE_OP (op ) \
131
153
template <class O , typename = std::enable_if_t <!detail::is_ann_ref_v<O>>> \
@@ -376,12 +398,10 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
376
398
377
399
operator T *() const noexcept = delete ;
378
400
379
- // T *get() const noexcept {
380
401
T *get () const noexcept {
381
402
#ifdef __SYCL_DEVICE_ONLY__
382
- return __builtin_intel_sycl_ptr_annotation (
383
- m_Ptr, detail::PropertyMetaInfo<Props>::name...,
384
- detail::PropertyMetaInfo<Props>::value...);
403
+ return annotationHelper<T, detail::annotation_filter<Props...>>::annotate (
404
+ m_Ptr);
385
405
#else
386
406
return m_Ptr;
387
407
#endif
0 commit comments