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