Skip to content

Commit e99f298

Browse files
author
Pavel Samolysov
authored
[SYCL] Add basic support for the generic_space address space (#5148)
The basic support is added to let the sycl::atomic_ref class template be instantiated with the address_space::generic_space address space.
1 parent 8ef9cee commit e99f298

File tree

6 files changed

+63
-35
lines changed

6 files changed

+63
-35
lines changed

‎sycl/include/CL/__spirv/spirv_ops.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -255,7 +255,7 @@ extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType,
255255

256256
#define __SPIRV_ATOMICS(macro, Arg) \
257257
macro(__attribute__((opencl_global)), Arg) \
258-
macro(__attribute__((opencl_local)), Arg)
258+
macro(__attribute__((opencl_local)), Arg) macro(, Arg)
259259

260260
__SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, float)
261261
__SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, double)

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -162,6 +162,11 @@ struct DecoratedType<ElementType, access::address_space::private_space> {
162162
using type = __OPENCL_PRIVATE_AS__ ElementType;
163163
};
164164

165+
template <typename ElementType>
166+
struct DecoratedType<ElementType, access::address_space::generic_space> {
167+
using type = ElementType;
168+
};
169+
165170
template <typename ElementType>
166171
struct DecoratedType<ElementType, access::address_space::global_space> {
167172
using type = __OPENCL_GLOBAL_AS__ ElementType;

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

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -114,11 +114,6 @@ template <> struct bit_equal<double> {
114114
template <typename T, memory_order DefaultOrder, memory_scope DefaultScope,
115115
access::address_space AddressSpace>
116116
class atomic_ref_base {
117-
static_assert(
118-
AddressSpace != access::address_space::generic_space,
119-
"access::address_space::generic_space is a valid address space but the "
120-
"address space is not supported yet.");
121-
122117
static_assert(
123118
detail::IsValidAtomicRefType<T>::value,
124119
"Invalid atomic type. Valid types are int, unsigned int, long, "

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

Lines changed: 52 additions & 27 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 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>>
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 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>>
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 &&
118-
(Space == access::address_space::global_space ||
132+
(Space == access::address_space::generic_space ||
133+
Space == access::address_space::global_space ||
119134
Space == access::address_space::global_device_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::generic_space ||
147+
Space == access::address_space::local_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 and
174+
// 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 &&
165-
(Space == access::address_space::global_space ||
181+
(Space == access::address_space::generic_space ||
182+
Space == access::address_space::global_space ||
166183
Space == access::address_space::global_device_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::generic_space ||
198+
Space == access::address_space::local_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 &&
381-
(Space == access::address_space::global_space ||
400+
(Space == access::address_space::generic_space ||
401+
Space == access::address_space::global_space ||
382402
Space == access::address_space::global_device_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::generic_space ||
414+
Space == access::address_space::local_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 &&
501-
(Space == access::address_space::global_space ||
523+
(Space == access::address_space::generic_space ||
524+
Space == access::address_space::global_space ||
502525
Space == access::address_space::global_device_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
509-
template <typename ElementType, int dimensions, access::mode Mode,
510-
typename PropertyListT, access::address_space _Space = Space,
511-
typename = typename detail::enable_if_t<
512-
_Space == Space && Space == access::address_space::local_space>>
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)>>
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
@@ -16,6 +16,10 @@ namespace sycl {
1616
template <typename ElementType, access::address_space Space> class multi_ptr;
1717
// Template specialization aliases for different pointer address spaces
1818

19+
template <typename ElementType>
20+
using generic_ptr =
21+
multi_ptr<ElementType, access::address_space::generic_space>;
22+
1923
template <typename ElementType>
2024
using global_ptr = multi_ptr<ElementType, access::address_space::global_space>;
2125

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -o %t.out -Xclang -verify-ignore-unexpected=note
2-
2+
// expected-no-diagnostics
33
#include <CL/sycl/atomic_ref.hpp>
44

55
struct A {};
@@ -20,6 +20,5 @@ int main() {
2020
A* p = &a;
2121
auto ref_p = sycl::atomic_ref<A *, sycl::memory_order_acq_rel,
2222
sycl::memory_scope_device>(p);
23-
// expected-error@CL/sycl/atomic_ref.hpp:* {{"access::address_space::generic_space is a valid address space but the address space is not supported yet."}}
2423
return 0;
2524
}

0 commit comments

Comments
 (0)