DPC++ Runtime
Runtime libraries for oneAPI DPC++
access.hpp
Go to the documentation of this file.
1 //==---------------- access.hpp --- SYCL access ----------------------------==//
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 #pragma once
9 
10 #include <CL/__spirv/spirv_ops.hpp>
11 #include <sycl/detail/common.hpp>
12 #include <sycl/detail/defines.hpp>
13 
14 namespace sycl {
16 namespace access {
17 
18 enum class target {
19  global_buffer __SYCL2020_DEPRECATED("use 'target::device' instead") = 2014,
20  constant_buffer __SYCL2020_DEPRECATED("use 'target::device' instead") = 2015,
21  local __SYCL2020_DEPRECATED("use `local_accessor` instead") = 2016,
22  image = 2017,
23  host_buffer __SYCL2020_DEPRECATED("use 'host_accessor' instead") = 2018,
24  host_image = 2019,
25  image_array = 2020,
26  host_task = 2021,
27  device = global_buffer,
28 };
29 
30 enum class mode {
31  read = 1024,
32  write = 1025,
33  read_write = 1026,
34  discard_write = 1027,
35  discard_read_write = 1028,
36  atomic = 1029
37 };
38 
39 enum class fence_space {
40  local_space = 0,
41  global_space = 1,
43 };
44 
45 enum class placeholder { false_t = 0, true_t = 1 };
46 
47 enum class address_space : int {
48  private_space = 0,
49  global_space = 1,
50  constant_space __SYCL2020_DEPRECATED("sycl::access::address_space::constant_"
51  "space is deprecated since SYCL 2020") =
52  2,
53  local_space = 3,
56  generic_space = 6, // TODO generic_space address space is not supported yet
57 };
58 
59 enum class decorated : int { no = 0, yes = 1, legacy = 2 };
60 } // namespace access
61 
62 using access::target;
64 
65 enum class image_target : unsigned int { device = 0, host_task = 1 };
66 
67 template <access_mode mode> struct mode_tag_t {
68  explicit mode_tag_t() = default;
69 };
70 
71 template <access_mode mode, target trgt> struct mode_target_tag_t {
72  explicit mode_target_tag_t() = default;
73 };
74 
78 inline constexpr mode_target_tag_t<access_mode::read, target::constant_buffer>
80 inline constexpr mode_target_tag_t<access_mode::read, target::host_task>
86 
87 namespace detail {
88 
90  return T == access::target::host_buffer || T == access::target::host_image;
91 }
92 
93 constexpr bool modeNeedsOldData(access::mode m) {
94  return m == access::mode::read || m == access::mode::write ||
95  m == access::mode::read_write || m == access::mode::atomic;
96 }
97 
98 constexpr bool modeWritesNewData(access::mode m) {
99  return m != access::mode::read;
100 }
101 
102 template <access::decorated Decorated> struct NegateDecorated;
103 template <> struct NegateDecorated<access::decorated::yes> {
104  static constexpr access::decorated value = access::decorated::no;
105 };
106 template <> struct NegateDecorated<access::decorated::no> {
107  static constexpr access::decorated value = access::decorated::yes;
108 };
109 
110 #ifdef __SYCL_DEVICE_ONLY__
111 #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global))
112 #ifdef __ENABLE_USM_ADDR_SPACE__
113 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device))
114 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host))
115 #else
116 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global))
117 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global))
118 #endif // __ENABLE_USM_ADDR_SPACE__
119 #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
120 #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
121 #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
122 #else
123 #define __OPENCL_GLOBAL_AS__
124 #define __OPENCL_GLOBAL_DEVICE_AS__
125 #define __OPENCL_GLOBAL_HOST_AS__
126 #define __OPENCL_LOCAL_AS__
127 #define __OPENCL_CONSTANT_AS__
128 #define __OPENCL_PRIVATE_AS__
129 #endif
130 
131 template <access::target accessTarget> struct TargetToAS {
132  constexpr static access::address_space AS =
133  access::address_space::global_space;
134 };
135 
136 #ifdef __ENABLE_USM_ADDR_SPACE__
137 template <> struct TargetToAS<access::target::device> {
138  constexpr static access::address_space AS =
139  access::address_space::ext_intel_global_device_space;
140 };
141 #endif // __ENABLE_USM_ADDR_SPACE__
142 
143 template <> struct TargetToAS<access::target::local> {
144  constexpr static access::address_space AS =
145  access::address_space::local_space;
146 };
147 
148 template <> struct TargetToAS<access::target::constant_buffer> {
149  constexpr static access::address_space AS =
150  access::address_space::constant_space;
151 };
152 
153 template <typename ElementType, access::address_space addressSpace>
155 
156 template <typename ElementType>
157 struct DecoratedType<ElementType, access::address_space::private_space> {
158  using type = __OPENCL_PRIVATE_AS__ ElementType;
159 };
160 
161 template <typename ElementType>
162 struct DecoratedType<ElementType, access::address_space::generic_space> {
163  using type = ElementType;
164 };
165 
166 template <typename ElementType>
167 struct DecoratedType<ElementType, access::address_space::global_space> {
168  using type = __OPENCL_GLOBAL_AS__ ElementType;
169 };
170 
171 template <typename ElementType>
172 struct DecoratedType<ElementType,
173  access::address_space::ext_intel_global_device_space> {
174  using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType;
175 };
176 
177 template <typename ElementType>
178 struct DecoratedType<ElementType,
179  access::address_space::ext_intel_global_host_space> {
180  using type = __OPENCL_GLOBAL_HOST_AS__ ElementType;
181 };
182 
183 template <typename ElementType>
184 struct DecoratedType<ElementType, access::address_space::constant_space> {
185  // Current implementation of address spaces handling leads to possibility
186  // of emitting incorrect (in terms of OpenCL) address space casts from
187  // constant to generic (and vise-versa). So, global address space is used
188  // here instead of constant to avoid incorrect address space casts in the
189  // produced device code.
190 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
191  using type = const __OPENCL_GLOBAL_AS__ ElementType;
192 #else
193  using type = __OPENCL_GLOBAL_AS__ ElementType;
194 #endif
195 };
196 
197 template <typename ElementType>
198 struct DecoratedType<ElementType, access::address_space::local_space> {
199  using type = __OPENCL_LOCAL_AS__ ElementType;
200 };
201 
202 #ifdef __SYCL_DEVICE_ONLY__
203 template <class T> struct deduce_AS_impl {
204  // Undecorated pointers are considered generic.
205  // TODO: This assumes that the implementation uses generic as default. If
206  // address space inference is used this may need to change.
207  static constexpr access::address_space value =
208  access::address_space::generic_space;
209 };
210 
211 #ifdef __ENABLE_USM_ADDR_SPACE__
212 template <class T> struct deduce_AS_impl<__OPENCL_GLOBAL_DEVICE_AS__ T> {
213  static constexpr access::address_space value =
214  access::address_space::ext_intel_global_device_space;
215 };
216 
217 template <class T> struct deduce_AS_impl<__OPENCL_GLOBAL_HOST_AS__ T> {
218  static constexpr access::address_space value =
219  access::address_space::ext_intel_global_host_space;
220 };
221 #endif // __ENABLE_USM_ADDR_SPACE__
222 
223 template <class T> struct deduce_AS_impl<__OPENCL_GLOBAL_AS__ T> {
224  static constexpr access::address_space value =
225  access::address_space::global_space;
226 };
227 
228 template <class T> struct deduce_AS_impl<__OPENCL_PRIVATE_AS__ T> {
229  static constexpr access::address_space value =
230  access::address_space::private_space;
231 };
232 
233 template <class T> struct deduce_AS_impl<__OPENCL_LOCAL_AS__ T> {
234  static constexpr access::address_space value =
235  access::address_space::local_space;
236 };
237 
238 template <class T> struct deduce_AS_impl<__OPENCL_CONSTANT_AS__ T> {
239  static constexpr access::address_space value =
240  access::address_space::constant_space;
241 };
242 
243 template <class T>
244 struct deduce_AS
245  : deduce_AS_impl<
246  std::remove_pointer_t<std::remove_reference_t<std::remove_cv_t<T>>>> {
247 };
248 #endif
249 
250 } // namespace detail
251 
252 template <typename T> struct remove_decoration {
253  using type = T;
254 };
255 
256 // Propagate through const qualifier.
257 template <typename T> struct remove_decoration<const T> {
258  using type = const typename remove_decoration<T>::type;
259 };
260 
261 // Propagate through pointer.
262 template <typename T> struct remove_decoration<T *> {
263  using type = typename remove_decoration<T>::type *;
264 };
265 
266 // Propagate through const qualified pointer.
267 template <typename T> struct remove_decoration<const T *> {
268  using type = const typename remove_decoration<T>::type *;
269 };
270 
271 // Propagate through reference.
272 template <typename T> struct remove_decoration<T &> {
273  using type = typename remove_decoration<T>::type &;
274 };
275 
276 // Propagate through const qualified reference.
277 template <typename T> struct remove_decoration<const T &> {
278  using type = const typename remove_decoration<T>::type &;
279 };
280 
281 #ifdef __SYCL_DEVICE_ONLY__
282 template <typename T> struct remove_decoration<__OPENCL_GLOBAL_AS__ T> {
283  using type = T;
284 };
285 
286 #ifdef __ENABLE_USM_ADDR_SPACE__
287 template <typename T> struct remove_decoration<__OPENCL_GLOBAL_DEVICE_AS__ T> {
288  using type = T;
289 };
290 
291 template <typename T> struct remove_decoration<__OPENCL_GLOBAL_HOST_AS__ T> {
292  using type = T;
293 };
294 
295 #endif // __ENABLE_USM_ADDR_SPACE__
296 
297 template <typename T> struct remove_decoration<__OPENCL_PRIVATE_AS__ T> {
298  using type = T;
299 };
300 
301 template <typename T> struct remove_decoration<__OPENCL_LOCAL_AS__ T> {
302  using type = T;
303 };
304 
305 template <typename T> struct remove_decoration<__OPENCL_CONSTANT_AS__ T> {
306  using type = T;
307 };
308 #endif // __SYCL_DEVICE_ONLY__
309 
310 template <typename T>
312 
313 namespace detail {
314 
315 // Helper function for selecting appropriate casts between address spaces.
316 template <typename ToT, typename FromT> inline ToT cast_AS(FromT from) {
317 #ifdef __SYCL_DEVICE_ONLY__
318  constexpr access::address_space ToAS = deduce_AS<ToT>::value;
319  constexpr access::address_space FromAS = deduce_AS<FromT>::value;
320  if constexpr (FromAS == access::address_space::generic_space) {
321 #if defined(__NVPTX__) || defined(__AMDGCN__)
322  // TODO: NVPTX and AMDGCN backends do not currently support the
323  // __spirv_GenericCastToPtrExplicit_* builtins, so to work around this
324  // we do C-style casting. This may produce warnings when targetting
325  // these backends.
326  return (ToT)from;
327 #else
328  using ToElemT = std::remove_pointer_t<remove_decoration_t<ToT>>;
329  if constexpr (ToAS == access::address_space::global_space)
330  return __SYCL_GenericCastToPtrExplicit_ToGlobal<ToElemT>(from);
331  else if constexpr (ToAS == access::address_space::local_space)
332  return __SYCL_GenericCastToPtrExplicit_ToLocal<ToElemT>(from);
333  else if constexpr (ToAS == access::address_space::private_space)
334  return __SYCL_GenericCastToPtrExplicit_ToPrivate<ToElemT>(from);
335 #ifdef __ENABLE_USM_ADDR_SPACE__
336  else if constexpr (ToAS == access::address_space::
337  ext_intel_global_device_space ||
338  ToAS ==
339  access::address_space::ext_intel_global_host_space)
340  // For extended address spaces we do not currently have a SPIR-V
341  // conversion function, so we do a C-style cast. This may produce
342  // warnings.
343  return (ToT)from;
344 #endif // __ENABLE_USM_ADDR_SPACE__
345  else
346  return reinterpret_cast<ToT>(from);
347 #endif // defined(__NVPTX__) || defined(__AMDGCN__)
348  } else
349 #ifdef __ENABLE_USM_ADDR_SPACE__
350  if constexpr (FromAS == access::address_space::global_space &&
351  (ToAS ==
352  access::address_space::ext_intel_global_device_space ||
353  ToAS ==
354  access::address_space::ext_intel_global_host_space)) {
355  // Casting from global address space to the global device and host address
356  // spaces is allowed.
357  return (ToT)from;
358  } else
359 #endif // __ENABLE_USM_ADDR_SPACE__
360 #endif // __SYCL_DEVICE_ONLY__
361  {
362  return reinterpret_cast<ToT>(from);
363  }
364 }
365 
366 } // namespace detail
367 
368 #undef __OPENCL_GLOBAL_AS__
369 #undef __OPENCL_GLOBAL_DEVICE_AS__
370 #undef __OPENCL_GLOBAL_HOST_AS__
371 #undef __OPENCL_LOCAL_AS__
372 #undef __OPENCL_CONSTANT_AS__
373 #undef __OPENCL_PRIVATE_AS__
374 
375 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
376 } // namespace sycl
sycl::_V1::access::address_space::generic_space
@ generic_space
spirv_ops.hpp
sycl::_V1::__SYCL2020_DEPRECATED
signed char __SYCL2020_DEPRECATED
Definition: aliases.hpp:96
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::ext_intel_global_device_space >::type
__OPENCL_GLOBAL_DEVICE_AS__ ElementType type
Definition: access.hpp:174
__OPENCL_GLOBAL_HOST_AS__
#define __OPENCL_GLOBAL_HOST_AS__
Definition: access.hpp:125
sycl::_V1::access::mode::discard_read_write
@ discard_read_write
sycl::_V1::image
Defines a shared image data.
Definition: image.hpp:30
sycl::_V1::access::mode
mode
Definition: access.hpp:30
sycl::_V1::detail::modeWritesNewData
constexpr bool modeWritesNewData(access::mode m)
Definition: access.hpp:98
T
sycl::_V1::mode_tag_t
Definition: access.hpp:67
sycl::_V1::remove_decoration::type
T type
Definition: access.hpp:253
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::constant_space >::type
__OPENCL_GLOBAL_AS__ ElementType type
Definition: access.hpp:193
sycl::_V1::access::placeholder::false_t
@ false_t
sycl::_V1::access::address_space::ext_intel_global_host_space
@ ext_intel_global_host_space
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::local_space >::type
__OPENCL_LOCAL_AS__ ElementType type
Definition: access.hpp:199
sycl::_V1::access::fence_space::local_space
@ local_space
sycl::_V1::access::target::host_image
@ host_image
sycl::_V1::access::mode::atomic
@ atomic
sycl::_V1::access::decorated::legacy
@ legacy
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::mode_target_tag_t
Definition: access.hpp:71
__OPENCL_LOCAL_AS__
#define __OPENCL_LOCAL_AS__
Definition: access.hpp:126
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::global_space >::type
__OPENCL_GLOBAL_AS__ ElementType type
Definition: access.hpp:168
sycl::_V1::access::fence_space
fence_space
Definition: access.hpp:39
sycl::_V1::write_only_host_task
constexpr mode_target_tag_t< access_mode::write, target::host_task > write_only_host_task
Definition: access.hpp:85
sycl::_V1::remove_decoration_t
typename remove_decoration< T >::type remove_decoration_t
Definition: access.hpp:311
sycl::_V1::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::private_space >::type
__OPENCL_PRIVATE_AS__ ElementType type
Definition: access.hpp:158
sycl::_V1::read_only
constexpr mode_tag_t< access_mode::read > read_only
Definition: access.hpp:75
__OPENCL_CONSTANT_AS__
#define __OPENCL_CONSTANT_AS__
Definition: access.hpp:127
sycl::_V1::access::placeholder
placeholder
Definition: access.hpp:45
sycl::_V1::remove_decoration
Definition: access.hpp:252
sycl::_V1::access::placeholder::true_t
@ true_t
sycl::_V1::access::target::host_task
@ host_task
sycl::_V1::access::address_space::private_space
@ private_space
sycl::_V1::remove_decoration< const T >::type
const typename remove_decoration< T >::type type
Definition: access.hpp:258
defines.hpp
common.hpp
sycl::_V1::detail::NegateDecorated
Definition: access.hpp:102
sycl::_V1::detail::TargetToAS
Definition: access.hpp:131
__OPENCL_GLOBAL_DEVICE_AS__
#define __OPENCL_GLOBAL_DEVICE_AS__
Definition: access.hpp:124
sycl::_V1::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:49
sycl::_V1::access::target
target
Definition: access.hpp:18
sycl::_V1::access::fence_space::global_and_local
@ global_and_local
sycl::_V1::read_write
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:76
sycl::_V1::read_only_host_task
constexpr mode_target_tag_t< access_mode::read, target::host_task > read_only_host_task
Definition: access.hpp:81
sycl::_V1::access::decorated
decorated
Definition: access.hpp:59
sycl::_V1::write_only
constexpr mode_tag_t< access_mode::write > write_only
Definition: access.hpp:77
sycl::_V1::read_write_host_task
constexpr mode_target_tag_t< access_mode::read_write, target::host_task > read_write_host_task
Definition: access.hpp:83
sycl::_V1::access::decorated::no
@ no
__OPENCL_PRIVATE_AS__
#define __OPENCL_PRIVATE_AS__
Definition: access.hpp:128
sycl::_V1::access::fence_space::global_space
@ global_space
sycl::_V1::detail::modeNeedsOldData
constexpr bool modeNeedsOldData(access::mode m)
Definition: access.hpp:93
sycl::_V1::access::decorated::yes
@ yes
sycl::_V1::remove_decoration< const T & >::type
const typename remove_decoration< T >::type & type
Definition: access.hpp:278
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::generic_space >::type
ElementType type
Definition: access.hpp:163
sycl::_V1::detail::cast_AS
ToT cast_AS(FromT from)
Definition: access.hpp:316
sycl::_V1::access::address_space::ext_intel_global_device_space
@ ext_intel_global_device_space
sycl::_V1::remove_decoration< const T * >::type
const typename remove_decoration< T >::type * type
Definition: access.hpp:268
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::ext_intel_global_host_space >::type
__OPENCL_GLOBAL_HOST_AS__ ElementType type
Definition: access.hpp:180
sycl::_V1::detail::DecoratedType
Definition: access.hpp:154
sycl::_V1::read_constant
constexpr mode_target_tag_t< access_mode::read, target::constant_buffer > read_constant
Definition: access.hpp:79
sycl::_V1::remove_decoration< T & >::type
typename remove_decoration< T >::type & type
Definition: access.hpp:273
sycl::_V1::detail::isTargetHostAccess
constexpr bool isTargetHostAccess(access::target T)
Definition: access.hpp:89
__OPENCL_GLOBAL_AS__
#define __OPENCL_GLOBAL_AS__
Definition: access.hpp:123
sycl::_V1::access::mode::read
@ read
sycl::_V1::image_target
image_target
Definition: access.hpp:65
sycl::_V1::access::mode::discard_write
@ discard_write
sycl::_V1::remove_decoration< T * >::type
typename remove_decoration< T >::type * type
Definition: access.hpp:263
sycl::_V1::access::address_space
address_space
Definition: access.hpp:47
sycl::_V1::access::target::image_array
@ image_array
Definition: accessor.hpp:3158