Skip to content

Commit e7f9906

Browse files
author
Pavel Samolysov
committed
[SYCL] Fix compilation errors when address_space::generic_space is used
1 parent c83ce4b commit e7f9906

File tree

5 files changed

+76
-29
lines changed

5 files changed

+76
-29
lines changed

‎sycl/include/CL/sycl/access/access.hpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -123,13 +123,15 @@ constexpr bool modeWritesNewData(access::mode m) {
123123
#define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
124124
#define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
125125
#define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
126+
#define __OPENCL_GENERIC_AS__ __attribute__((opencl_generic))
126127
#else
127128
#define __OPENCL_GLOBAL_AS__
128129
#define __OPENCL_GLOBAL_DEVICE_AS__
129130
#define __OPENCL_GLOBAL_HOST_AS__
130131
#define __OPENCL_LOCAL_AS__
131132
#define __OPENCL_CONSTANT_AS__
132133
#define __OPENCL_PRIVATE_AS__
134+
#define __OPENCL_GENERIC_AS__
133135
#endif
134136

135137
template <access::target accessTarget> struct TargetToAS {
@@ -195,6 +197,12 @@ template <typename ElementType>
195197
struct DecoratedType<ElementType, access::address_space::local_space> {
196198
using type = __OPENCL_LOCAL_AS__ ElementType;
197199
};
200+
201+
template <typename ElementType>
202+
struct DecoratedType<ElementType, access::address_space::generic_space> {
203+
using type = __OPENCL_GENERIC_AS__ ElementType;
204+
};
205+
198206
template <class T> struct remove_AS { typedef T type; };
199207

200208
#ifdef __SYCL_DEVICE_ONLY__
@@ -261,6 +269,7 @@ template <class T> struct deduce_AS<__OPENCL_CONSTANT_AS__ T> {
261269
#undef __OPENCL_LOCAL_AS__
262270
#undef __OPENCL_CONSTANT_AS__
263271
#undef __OPENCL_PRIVATE_AS__
272+
#undef __OPENCL_GENERIC_AS__
264273
} // namespace detail
265274

266275
} // namespace sycl

‎sycl/include/CL/sycl/atomic_ref.hpp

Lines changed: 13 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -39,6 +39,15 @@ template <typename T> struct IsValidAtomicRefType {
3939
std::is_pointer<T>::value);
4040
};
4141

42+
template <cl::sycl::access::address_space AS>
43+
struct IsValidAtomicRefAddressSpace {
44+
static constexpr bool value =
45+
(AS == access::address_space::global_space ||
46+
AS == access::address_space::local_space ||
47+
AS == access::address_space::global_device_space ||
48+
AS == access::address_space::generic_space);
49+
};
50+
4251
// DefaultOrder parameter is limited to read-modify-write orders
4352
template <memory_order Order>
4453
using IsValidDefaultOrder = bool_constant<Order == memory_order::relaxed ||
@@ -110,9 +119,10 @@ class atomic_ref_base {
110119
"Invalid atomic type. Valid types are int, unsigned int, long, "
111120
"unsigned long, long long, unsigned long long, float, double "
112121
"and pointer types");
113-
static_assert(detail::IsValidAtomicAddressSpace<AddressSpace>::value,
114-
"Invalid atomic address_space. Valid address spaces are: "
115-
"global_space, local_space, global_device_space");
122+
static_assert(
123+
detail::IsValidAtomicRefAddressSpace<AddressSpace>::value,
124+
"Invalid atomic address_space. Valid address spaces are: "
125+
"global_space, local_space, global_device_space, generic_space");
116126
static_assert(
117127
detail::IsValidDefaultOrder<DefaultOrder>::value,
118128
"Invalid default memory_order for atomics. Valid defaults are: "

‎sycl/include/CL/sycl/multi_ptr.hpp

Lines changed: 49 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,13 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
5151
multi_ptr(const multi_ptr &rhs) = default;
5252
multi_ptr(multi_ptr &&) = default;
5353
#ifdef __SYCL_DEVICE_ONLY__
54+
// The attribute opencl_generic used for the generic address space generates
55+
// nothing and 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>>
5461
multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
5562
#endif
5663

@@ -71,6 +78,13 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
7178
multi_ptr &operator=(multi_ptr &&) = default;
7279

7380
#ifdef __SYCL_DEVICE_ONLY__
81+
// The attribute opencl_generic used for the generic address space generates
82+
// nothing and 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>>
7488
multi_ptr &operator=(pointer_t pointer) {
7589
m_Pointer = pointer;
7690
return *this;
@@ -109,26 +123,28 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
109123
return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
110124
}
111125

112-
// Only if Space == global_space || global_device_space
126+
// Only if Space == global_space || global_device_space || generic_space
113127
template <int dimensions, access::mode Mode,
114128
access::placeholder isPlaceholder, typename PropertyListT,
115129
access::address_space _Space = Space,
116130
typename = typename detail::enable_if_t<
117131
_Space == Space &&
118132
(Space == access::address_space::global_space ||
119-
Space == access::address_space::global_device_space)>>
133+
Space == access::address_space::global_device_space ||
134+
Space == access::address_space::generic_space)>>
120135
multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
121136
isPlaceholder, PropertyListT>
122137
Accessor) {
123138
m_Pointer = (pointer_t)(Accessor.get_pointer().get());
124139
}
125140

126-
// Only if Space == local_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 && Space == access::address_space::local_space>>
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::local_space ||
147+
Space == access::address_space::generic_space)>>
132148
multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
133149
isPlaceholder, PropertyListT>
134150
Accessor)
@@ -154,29 +170,32 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
154170
// 2. from multi_ptr<ElementType, Space> to multi_ptr<const ElementType,
155171
// Space>
156172

157-
// Only if Space == global_space || global_device_space and element type is
158-
// const
173+
// Only if Space == global_space || global_device_space || generic_space
174+
// and element type is const
159175
template <
160176
int dimensions, access::mode Mode, access::placeholder isPlaceholder,
161177
typename PropertyListT, access::address_space _Space = Space,
162178
typename ET = ElementType,
163179
typename = typename detail::enable_if_t<
164180
_Space == Space &&
165181
(Space == access::address_space::global_space ||
166-
Space == access::address_space::global_device_space) &&
182+
Space == access::address_space::global_device_space ||
183+
Space == access::address_space::generic_space) &&
167184
std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
168185
multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
169186
access::target::device, isPlaceholder, PropertyListT>
170187
Accessor)
171188
: multi_ptr(Accessor.get_pointer()) {}
172189

173-
// Only if Space == local_space and element type is const
190+
// Only if Space == local_space || generic_space and element type is const
174191
template <
175192
int dimensions, access::mode Mode, access::placeholder isPlaceholder,
176193
typename PropertyListT, access::address_space _Space = Space,
177194
typename ET = ElementType,
178195
typename = typename detail::enable_if_t<
179-
_Space == Space && Space == access::address_space::local_space &&
196+
_Space == Space &&
197+
(Space == access::address_space::local_space ||
198+
Space == access::address_space::generic_space) &&
180199
std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
181200
multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
182201
access::target::local, isPlaceholder, PropertyListT>
@@ -373,23 +392,26 @@ template <access::address_space Space> class multi_ptr<void, Space> {
373392
return *this;
374393
}
375394

376-
// Only if Space == global_space || global_device_space
395+
// Only if Space == global_space || global_device_space || generic_space
377396
template <typename ElementType, int dimensions, access::mode Mode,
378397
typename PropertyListT, access::address_space _Space = Space,
379398
typename = typename detail::enable_if_t<
380399
_Space == Space &&
381400
(Space == access::address_space::global_space ||
382-
Space == access::address_space::global_device_space)>>
401+
Space == access::address_space::global_device_space ||
402+
Space == access::address_space::generic_space)>>
383403
multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
384404
access::placeholder::false_t, PropertyListT>
385405
Accessor)
386406
: multi_ptr(Accessor.get_pointer()) {}
387407

388-
// Only if Space == local_space
389-
template <typename ElementType, int dimensions, access::mode Mode,
390-
typename PropertyListT, access::address_space _Space = Space,
391-
typename = typename detail::enable_if_t<
392-
_Space == Space && Space == access::address_space::local_space>>
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::local_space ||
414+
Space == access::address_space::generic_space)>>
393415
multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
394416
access::placeholder::false_t, PropertyListT>
395417
Accessor)
@@ -493,23 +515,26 @@ class multi_ptr<const void, Space> {
493515
return *this;
494516
}
495517

496-
// Only if Space == global_space || global_device_space
518+
// Only if Space == global_space || global_device_space || generic_space
497519
template <typename ElementType, int dimensions, access::mode Mode,
498520
typename PropertyListT, access::address_space _Space = Space,
499521
typename = typename detail::enable_if_t<
500522
_Space == Space &&
501523
(Space == access::address_space::global_space ||
502-
Space == access::address_space::global_device_space)>>
524+
Space == access::address_space::global_device_space ||
525+
Space == access::address_space::generic_space)>>
503526
multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
504527
access::placeholder::false_t, PropertyListT>
505528
Accessor)
506529
: multi_ptr(Accessor.get_pointer()) {}
507530

508-
// Only if Space == local_space
531+
// Only if Space == local_space || generic_space
509532
template <typename ElementType, int dimensions, access::mode Mode,
510533
typename PropertyListT, access::address_space _Space = Space,
511534
typename = typename detail::enable_if_t<
512-
_Space == Space && Space == access::address_space::local_space>>
535+
_Space == Space &&
536+
(Space == access::address_space::local_space ||
537+
Space == access::address_space::generic_space)>>
513538
multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
514539
access::placeholder::false_t, PropertyListT>
515540
Accessor)

‎sycl/include/CL/sycl/pointers.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,5 +38,9 @@ template <typename ElementType>
3838
using private_ptr =
3939
multi_ptr<ElementType, access::address_space::private_space>;
4040

41+
template <typename ElementType>
42+
using generic_ptr =
43+
multi_ptr<ElementType, access::address_space::generic_space>;
44+
4145
} // namespace sycl
4246
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/test/basic_tests/atomic-ref-instantiation.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,8 +21,7 @@ int main() {
2121
A a;
2222
A* p = &a;
2323
auto ref_p = sycl::atomic_ref<A *, sycl::memory_order_acq_rel,
24-
sycl::memory_scope_device,
25-
sycl::access::address_space::local_space>(p);
24+
sycl::memory_scope_device>(p);
2625

2726
return 0;
2827
}

0 commit comments

Comments
 (0)