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 <sycl/detail/common.hpp>
11 #include <sycl/detail/defines.hpp>
12 
13 namespace sycl {
15 namespace access {
16 
17 enum class target {
18  global_buffer __SYCL2020_DEPRECATED("use 'target::device' instead") = 2014,
19  constant_buffer = 2015,
20  local __SYCL2020_DEPRECATED("use `local_accessor` instead") = 2016,
21  image = 2017,
22  host_buffer = 2018,
23  host_image = 2019,
24  image_array = 2020,
25  device = global_buffer,
26 };
27 
28 enum class mode {
29  read = 1024,
30  write = 1025,
31  read_write = 1026,
32  discard_write = 1027,
33  discard_read_write = 1028,
34  atomic = 1029
35 };
36 
37 enum class fence_space {
38  local_space = 0,
39  global_space = 1,
41 };
42 
43 enum class placeholder { false_t = 0, true_t = 1 };
44 
45 enum class address_space : int {
46  private_space = 0,
47  global_space = 1,
48  constant_space __SYCL2020_DEPRECATED("sycl::access::address_space::constant_"
49  "space is deprecated since SYCL 2020") =
50  2,
51  local_space = 3,
54  generic_space = 6, // TODO generic_space address space is not supported yet
55 };
56 
57 enum class decorated : int { no = 0, yes = 1, legacy = 2 };
58 } // namespace access
59 
60 using access::target;
62 
63 template <access_mode mode> struct mode_tag_t {
64  explicit mode_tag_t() = default;
65 };
66 
67 template <access_mode mode, target trgt> struct mode_target_tag_t {
68  explicit mode_target_tag_t() = default;
69 };
70 
71 #if __cplusplus >= 201703L
72 
73 inline constexpr mode_tag_t<access_mode::read> read_only{};
74 inline constexpr mode_tag_t<access_mode::read_write> read_write{};
75 inline constexpr mode_tag_t<access_mode::write> write_only{};
76 inline constexpr mode_target_tag_t<access_mode::read, target::constant_buffer>
77  read_constant{};
78 
79 #else
80 
81 namespace {
82 
83 constexpr const auto &read_only =
84  sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::read>>::value;
85 constexpr const auto &read_write = sycl::detail::InlineVariableHelper<
86  mode_tag_t<access_mode::read_write>>::value;
87 constexpr const auto &write_only =
88  sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::write>>::value;
89 constexpr const auto &read_constant = sycl::detail::InlineVariableHelper<
90  mode_target_tag_t<access_mode::read, target::constant_buffer>>::value;
91 
92 } // namespace
93 
94 #endif
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 ||
104  m == access::mode::read_write || m == access::mode::atomic;
105 }
106 
107 constexpr bool modeWritesNewData(access::mode m) {
108  return m != access::mode::read;
109 }
110 
111 #ifdef __SYCL_DEVICE_ONLY__
112 #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global))
113 #ifdef __ENABLE_USM_ADDR_SPACE__
114 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device))
115 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host))
116 #else
117 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global))
118 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global))
119 #endif // __ENABLE_USM_ADDR_SPACE__
120 #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
121 #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
122 #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
123 #else
124 #define __OPENCL_GLOBAL_AS__
125 #define __OPENCL_GLOBAL_DEVICE_AS__
126 #define __OPENCL_GLOBAL_HOST_AS__
127 #define __OPENCL_LOCAL_AS__
128 #define __OPENCL_CONSTANT_AS__
129 #define __OPENCL_PRIVATE_AS__
130 #endif
131 
132 template <access::target accessTarget> struct TargetToAS {
133  constexpr static access::address_space AS =
134  access::address_space::global_space;
135 };
136 
137 #ifdef __ENABLE_USM_ADDR_SPACE__
138 template <> struct TargetToAS<access::target::device> {
139  constexpr static access::address_space AS =
140  access::address_space::ext_intel_global_device_space;
141 };
142 #endif // __ENABLE_USM_ADDR_SPACE__
143 
144 template <> struct TargetToAS<access::target::local> {
145  constexpr static access::address_space AS =
146  access::address_space::local_space;
147 };
148 
149 template <> struct TargetToAS<access::target::constant_buffer> {
150  constexpr static access::address_space AS =
151  access::address_space::constant_space;
152 };
153 
154 template <typename ElementType, access::address_space addressSpace>
156 
157 template <typename ElementType>
158 struct DecoratedType<ElementType, access::address_space::private_space> {
159  using type = __OPENCL_PRIVATE_AS__ ElementType;
160 };
161 
162 template <typename ElementType>
163 struct DecoratedType<ElementType, access::address_space::generic_space> {
164  using type = ElementType;
165 };
166 
167 template <typename ElementType>
168 struct DecoratedType<ElementType, access::address_space::global_space> {
169  using type = __OPENCL_GLOBAL_AS__ ElementType;
170 };
171 
172 template <typename ElementType>
173 struct DecoratedType<ElementType,
174  access::address_space::ext_intel_global_device_space> {
175  using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType;
176 };
177 
178 template <typename ElementType>
179 struct DecoratedType<ElementType,
180  access::address_space::ext_intel_global_host_space> {
181  using type = __OPENCL_GLOBAL_HOST_AS__ ElementType;
182 };
183 
184 template <typename ElementType>
185 struct DecoratedType<ElementType, access::address_space::constant_space> {
186  // Current implementation of address spaces handling leads to possibility
187  // of emitting incorrect (in terms of OpenCL) address space casts from
188  // constant to generic (and vise-versa). So, global address space is used
189  // here instead of constant to avoid incorrect address space casts in the
190  // produced device code.
191 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
192  using type = const __OPENCL_GLOBAL_AS__ ElementType;
193 #else
194  using type = __OPENCL_GLOBAL_AS__ ElementType;
195 #endif
196 };
197 
198 template <typename ElementType>
199 struct DecoratedType<ElementType, access::address_space::local_space> {
200  using type = __OPENCL_LOCAL_AS__ ElementType;
201 };
202 template <class T> struct remove_AS {
203  typedef T type;
204 };
205 
206 #ifdef __SYCL_DEVICE_ONLY__
207 template <class T> struct deduce_AS {
208  // Undecorated pointers are considered generic.
209  // TODO: This assumes that the implementation uses generic as default. If
210  // address space inference is used this may need to change.
211  static const access::address_space value =
212  access::address_space::generic_space;
213 };
214 
215 template <class T> struct remove_AS<__OPENCL_GLOBAL_AS__ T> {
216  typedef T type;
217 };
218 
219 #ifdef __ENABLE_USM_ADDR_SPACE__
220 template <class T> struct remove_AS<__OPENCL_GLOBAL_DEVICE_AS__ T> {
221  typedef T type;
222 };
223 
224 template <class T> struct remove_AS<__OPENCL_GLOBAL_HOST_AS__ T> {
225  typedef T type;
226 };
227 
228 template <class T> struct deduce_AS<__OPENCL_GLOBAL_DEVICE_AS__ T> {
229  static const access::address_space value =
230  access::address_space::ext_intel_global_device_space;
231 };
232 
233 template <class T> struct deduce_AS<__OPENCL_GLOBAL_HOST_AS__ T> {
234  static const access::address_space value =
235  access::address_space::ext_intel_global_host_space;
236 };
237 #endif // __ENABLE_USM_ADDR_SPACE__
238 
239 template <class T> struct remove_AS<__OPENCL_PRIVATE_AS__ T> {
240  typedef T type;
241 };
242 
243 template <class T> struct remove_AS<__OPENCL_LOCAL_AS__ T> {
244  typedef T type;
245 };
246 
247 template <class T> struct remove_AS<__OPENCL_CONSTANT_AS__ T> {
248  typedef T type;
249 };
250 
251 template <class T> struct deduce_AS<__OPENCL_GLOBAL_AS__ T> {
252  static const access::address_space value =
253  access::address_space::global_space;
254 };
255 
256 template <class T> struct deduce_AS<__OPENCL_PRIVATE_AS__ T> {
257  static const access::address_space value =
258  access::address_space::private_space;
259 };
260 
261 template <class T> struct deduce_AS<__OPENCL_LOCAL_AS__ T> {
262  static const access::address_space value = access::address_space::local_space;
263 };
264 
265 template <class T> struct deduce_AS<__OPENCL_CONSTANT_AS__ T> {
266  static const access::address_space value =
267  access::address_space::constant_space;
268 };
269 #endif
270 
271 #undef __OPENCL_GLOBAL_AS__
272 #undef __OPENCL_GLOBAL_DEVICE_AS__
273 #undef __OPENCL_GLOBAL_HOST_AS__
274 #undef __OPENCL_LOCAL_AS__
275 #undef __OPENCL_CONSTANT_AS__
276 #undef __OPENCL_PRIVATE_AS__
277 } // namespace detail
278 
279 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
280 } // namespace sycl
sycl::_V1::access::address_space::generic_space
@ generic_space
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::ext_intel_global_device_space >::type
__OPENCL_GLOBAL_DEVICE_AS__ ElementType type
Definition: access.hpp:175
__OPENCL_GLOBAL_HOST_AS__
#define __OPENCL_GLOBAL_HOST_AS__
Definition: access.hpp:126
sycl::_V1::access::mode::discard_read_write
@ discard_read_write
sycl::_V1::image
Defines a shared image data.
Definition: image.hpp:181
sycl::_V1::access::mode
mode
Definition: access.hpp:28
sycl::_V1::detail::modeWritesNewData
constexpr bool modeWritesNewData(access::mode m)
Definition: access.hpp:107
sycl::_V1::mode_tag_t
Definition: access.hpp:63
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::constant_space >::type
__OPENCL_GLOBAL_AS__ ElementType type
Definition: access.hpp:194
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:13
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::local_space >::type
__OPENCL_LOCAL_AS__ ElementType type
Definition: access.hpp:200
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:13
sycl::_V1::mode_target_tag_t
Definition: access.hpp:67
sycl::_V1::info::global_mem_cache_type::read_only
@ read_only
__OPENCL_LOCAL_AS__
#define __OPENCL_LOCAL_AS__
Definition: access.hpp:127
sycl::_V1::access::target::host_buffer
@ host_buffer
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::global_space >::type
__OPENCL_GLOBAL_AS__ ElementType type
Definition: access.hpp:169
sycl::_V1::access::fence_space
fence_space
Definition: access.hpp:37
sycl::_V1::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:109
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::private_space >::type
__OPENCL_PRIVATE_AS__ ElementType type
Definition: access.hpp:159
__OPENCL_CONSTANT_AS__
#define __OPENCL_CONSTANT_AS__
Definition: access.hpp:128
sycl::_V1::access::placeholder
placeholder
Definition: access.hpp:43
sycl::_V1::access::placeholder::true_t
@ true_t
sycl::_V1::access::address_space::private_space
@ private_space
defines.hpp
sycl::_V1::detail::remove_AS::type
T type
Definition: access.hpp:203
common.hpp
sycl::_V1::detail::TargetToAS
Definition: access.hpp:132
__OPENCL_GLOBAL_DEVICE_AS__
#define __OPENCL_GLOBAL_DEVICE_AS__
Definition: access.hpp:125
sycl::_V1::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:47
sycl::_V1::access::target
target
Definition: access.hpp:17
sycl::_V1::access::fence_space::global_and_local
@ global_and_local
sycl::_V1::access::decorated
decorated
Definition: access.hpp:57
sycl::_V1::access::decorated::no
@ no
__OPENCL_PRIVATE_AS__
#define __OPENCL_PRIVATE_AS__
Definition: access.hpp:129
sycl::_V1::access::fence_space::global_space
@ global_space
sycl::_V1::detail::remove_AS
Definition: access.hpp:202
sycl::_V1::detail::modeNeedsOldData
constexpr bool modeNeedsOldData(access::mode m)
Definition: access.hpp:102
sycl::_V1::access::decorated::yes
@ yes
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::generic_space >::type
ElementType type
Definition: access.hpp:164
sycl::_V1::access::address_space::ext_intel_global_device_space
@ ext_intel_global_device_space
sycl::_V1::detail::DecoratedType< ElementType, access::address_space::ext_intel_global_host_space >::type
__OPENCL_GLOBAL_HOST_AS__ ElementType type
Definition: access.hpp:181
sycl::_V1::detail::DecoratedType
Definition: access.hpp:155
sycl::_V1::access::target::constant_buffer
@ constant_buffer
sycl::_V1::access::mode::read_write
@ read_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:124
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:54
sycl::_V1::access::mode::read
@ read
sycl::_V1::access::mode::discard_write
@ discard_write
sycl::_V1::access::address_space
address_space
Definition: access.hpp:45
sycl::_V1::access::target::image_array
@ image_array
Definition: accessor.hpp:2776