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