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 
12 
14 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 = 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 = 2,
49  local_space = 3,
52  global_device_space __SYCL2020_DEPRECATED(
53  "use 'ext_intel_global_device_space' instead") =
55  global_host_space __SYCL2020_DEPRECATED(
56  "use 'ext_intel_global_host_space' instead") =
58  ext_intel_host_device_space __SYCL2020_DEPRECATED(
59  "use 'ext_intel_global_host_space' instead") =
61  generic_space = 6, // TODO generic_space address space is not supported yet
62 };
63 
64 enum class decorated : int { no = 0, yes = 1, legacy = 2 };
65 } // namespace access
66 
67 using access::target;
69 
70 template <access_mode mode> struct mode_tag_t {
71  explicit mode_tag_t() = default;
72 };
73 
74 template <access_mode mode, target trgt> struct mode_target_tag_t {
75  explicit mode_target_tag_t() = default;
76 };
77 
78 #if __cplusplus >= 201703L
79 
80 inline constexpr mode_tag_t<access_mode::read> read_only{};
81 inline constexpr mode_tag_t<access_mode::read_write> read_write{};
82 inline constexpr mode_tag_t<access_mode::write> write_only{};
83 inline constexpr mode_target_tag_t<access_mode::read, target::constant_buffer>
84  read_constant{};
85 
86 #else
87 
88 namespace {
89 
90 constexpr const auto &read_only =
91  sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::read>>::value;
92 constexpr const auto &read_write = sycl::detail::InlineVariableHelper<
93  mode_tag_t<access_mode::read_write>>::value;
94 constexpr const auto &write_only =
95  sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::write>>::value;
96 constexpr const auto &read_constant = sycl::detail::InlineVariableHelper<
97  mode_target_tag_t<access_mode::read, target::constant_buffer>>::value;
98 
99 } // namespace
100 
101 #endif
102 
103 namespace detail {
104 
106  return T == access::target::host_buffer || T == access::target::host_image;
107 }
108 
109 constexpr bool modeNeedsOldData(access::mode m) {
110  return m == access::mode::read || m == access::mode::write ||
111  m == access::mode::read_write || m == access::mode::atomic;
112 }
113 
114 constexpr bool modeWritesNewData(access::mode m) {
115  return m != access::mode::read;
116 }
117 
118 #ifdef __SYCL_DEVICE_ONLY__
119 #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global))
120 #ifdef __ENABLE_USM_ADDR_SPACE__
121 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device))
122 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host))
123 #else
124 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global))
125 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global))
126 #endif // __ENABLE_USM_ADDR_SPACE__
127 #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
128 #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
129 #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
130 #else
131 #define __OPENCL_GLOBAL_AS__
132 #define __OPENCL_GLOBAL_DEVICE_AS__
133 #define __OPENCL_GLOBAL_HOST_AS__
134 #define __OPENCL_LOCAL_AS__
135 #define __OPENCL_CONSTANT_AS__
136 #define __OPENCL_PRIVATE_AS__
137 #endif
138 
139 template <access::target accessTarget> struct TargetToAS {
140  constexpr static access::address_space AS =
141  access::address_space::global_space;
142 };
143 
144 #ifdef __ENABLE_USM_ADDR_SPACE__
145 template <> struct TargetToAS<access::target::device> {
146  constexpr static access::address_space AS =
147  access::address_space::ext_intel_global_device_space;
148 };
149 #endif // __ENABLE_USM_ADDR_SPACE__
150 
151 template <> struct TargetToAS<access::target::local> {
152  constexpr static access::address_space AS =
153  access::address_space::local_space;
154 };
155 
156 template <> struct TargetToAS<access::target::constant_buffer> {
157  constexpr static access::address_space AS =
158  access::address_space::constant_space;
159 };
160 
161 template <typename ElementType, access::address_space addressSpace>
163 
164 template <typename ElementType>
165 struct DecoratedType<ElementType, access::address_space::private_space> {
166  using type = __OPENCL_PRIVATE_AS__ ElementType;
167 };
168 
169 template <typename ElementType>
170 struct DecoratedType<ElementType, access::address_space::generic_space> {
171  using type = ElementType;
172 };
173 
174 template <typename ElementType>
175 struct DecoratedType<ElementType, access::address_space::global_space> {
176  using type = __OPENCL_GLOBAL_AS__ ElementType;
177 };
178 
179 template <typename ElementType>
180 struct DecoratedType<ElementType,
181  access::address_space::ext_intel_global_device_space> {
182  using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType;
183 };
184 
185 template <typename ElementType>
186 struct DecoratedType<ElementType,
187  access::address_space::ext_intel_global_host_space> {
188  using type = __OPENCL_GLOBAL_HOST_AS__ ElementType;
189 };
190 
191 template <typename ElementType>
192 struct DecoratedType<ElementType, access::address_space::constant_space> {
193  // Current implementation of address spaces handling leads to possibility
194  // of emitting incorrect (in terms of OpenCL) address space casts from
195  // constant to generic (and vise-versa). So, global address space is used here
196  // instead of constant to avoid incorrect address space casts in the produced
197  // device code.
198 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
199  using type = const __OPENCL_GLOBAL_AS__ ElementType;
200 #else
201  using type = __OPENCL_GLOBAL_AS__ ElementType;
202 #endif
203 };
204 
205 template <typename ElementType>
206 struct DecoratedType<ElementType, access::address_space::local_space> {
207  using type = __OPENCL_LOCAL_AS__ ElementType;
208 };
209 template <class T> struct remove_AS { typedef T type; };
210 
211 #ifdef __SYCL_DEVICE_ONLY__
212 template <class T> struct deduce_AS {
213  static_assert(!std::is_same<typename detail::remove_AS<T>::type, T>::value,
214  "Only types with address space attributes are supported");
215 };
216 
217 template <class T> struct remove_AS<__OPENCL_GLOBAL_AS__ T> { typedef T type; };
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> { typedef T type; };
244 
245 template <class T> struct remove_AS<__OPENCL_CONSTANT_AS__ T> {
246  typedef T type;
247 };
248 
249 template <class T> struct deduce_AS<__OPENCL_GLOBAL_AS__ T> {
250  static const access::address_space value =
251  access::address_space::global_space;
252 };
253 
254 template <class T> struct deduce_AS<__OPENCL_PRIVATE_AS__ T> {
255  static const access::address_space value =
256  access::address_space::private_space;
257 };
258 
259 template <class T> struct deduce_AS<__OPENCL_LOCAL_AS__ T> {
260  static const access::address_space value = access::address_space::local_space;
261 };
262 
263 template <class T> struct deduce_AS<__OPENCL_CONSTANT_AS__ T> {
264  static const access::address_space value =
265  access::address_space::constant_space;
266 };
267 #endif
268 
269 #undef __OPENCL_GLOBAL_AS__
270 #undef __OPENCL_GLOBAL_DEVICE_AS__
271 #undef __OPENCL_GLOBAL_HOST_AS__
272 #undef __OPENCL_LOCAL_AS__
273 #undef __OPENCL_CONSTANT_AS__
274 #undef __OPENCL_PRIVATE_AS__
275 } // namespace detail
276 
277 } // namespace sycl
278 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::DecoratedType< ElementType, access::address_space::generic_space >::type
ElementType type
Definition: access.hpp:171
__OPENCL_GLOBAL_HOST_AS__
#define __OPENCL_GLOBAL_HOST_AS__
Definition: access.hpp:133
T
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:56
cl::sycl::access::placeholder
placeholder
Definition: access.hpp:43
cl::sycl::access::mode::read
@ read
cl::sycl::detail::DecoratedType< ElementType, access::address_space::ext_intel_global_device_space >::type
__OPENCL_GLOBAL_DEVICE_AS__ ElementType type
Definition: access.hpp:182
cl::sycl::access::target::local
@ local
cl::sycl::access::target::host_image
@ host_image
cl::sycl::access::target::image_array
@ image_array
cl::sycl::access::mode::atomic
@ atomic
sycl
Definition: invoke_simd.hpp:68
cl::sycl::access::address_space::constant_space
@ constant_space
__OPENCL_LOCAL_AS__
#define __OPENCL_LOCAL_AS__
Definition: access.hpp:134
cl::sycl::access::fence_space::local_space
@ local_space
cl::sycl::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
cl::sycl::detail::DecoratedType< ElementType, access::address_space::global_space >::type
__OPENCL_GLOBAL_AS__ ElementType type
Definition: access.hpp:176
cl::sycl::access::decorated::legacy
@ legacy
cl::sycl::access::placeholder::false_t
@ false_t
cl::sycl::access::address_space::generic_space
@ generic_space
cl::sycl::access::fence_space
fence_space
Definition: access.hpp:37
__OPENCL_CONSTANT_AS__
#define __OPENCL_CONSTANT_AS__
Definition: access.hpp:135
cl::sycl::access::decorated
decorated
Definition: access.hpp:64
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:36
cl::sycl::access::target::host_buffer
@ host_buffer
cl::sycl::detail::TargetToAS
Definition: access.hpp:139
defines.hpp
cl::sycl::access::target
target
Definition: access.hpp:17
cl::sycl::access::placeholder::true_t
@ true_t
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::access::decorated::no
@ no
__OPENCL_GLOBAL_DEVICE_AS__
#define __OPENCL_GLOBAL_DEVICE_AS__
Definition: access.hpp:132
cl::sycl::access::mode::discard_read_write
@ discard_read_write
cl::sycl::detail::DecoratedType< ElementType, access::address_space::private_space >::type
__OPENCL_PRIVATE_AS__ ElementType type
Definition: access.hpp:166
cl::sycl::detail::DecoratedType< ElementType, access::address_space::ext_intel_global_host_space >::type
__OPENCL_GLOBAL_HOST_AS__ ElementType type
Definition: access.hpp:188
cl::sycl::image
Defines a shared image data.
Definition: image_impl.hpp:29
cl::sycl::detail::DecoratedType< ElementType, access::address_space::constant_space >::type
__OPENCL_GLOBAL_AS__ ElementType type
Definition: access.hpp:201
cl::sycl::access::address_space
address_space
Definition: access.hpp:45
cl::sycl::access::decorated::yes
@ yes
cl::sycl::access::address_space::ext_intel_global_host_space
@ ext_intel_global_host_space
cl::sycl::access::target::constant_buffer
@ constant_buffer
cl::sycl::access::fence_space::global_space
@ global_space
cl::sycl::detail::remove_AS
Definition: access.hpp:209
cl::sycl::access::address_space::ext_intel_global_device_space
@ ext_intel_global_device_space
cl::sycl::detail::DecoratedType< ElementType, access::address_space::local_space >::type
__OPENCL_LOCAL_AS__ ElementType type
Definition: access.hpp:207
__OPENCL_PRIVATE_AS__
#define __OPENCL_PRIVATE_AS__
Definition: access.hpp:136
cl::sycl::access::mode::discard_write
@ discard_write
cl::sycl::mode_target_tag_t
Definition: access.hpp:74
cl::sycl::access::mode::read_write
@ read_write
common.hpp
cl::sycl::access::fence_space::global_and_local
@ global_and_local
cl::sycl::access::mode
mode
Definition: access.hpp:28
cl::sycl::mode_tag_t
Definition: access.hpp:70
cl::sycl::info::device
device
Definition: info_desc.hpp:53
cl::sycl::detail::modeWritesNewData
constexpr bool modeWritesNewData(access::mode m)
Definition: access.hpp:114
cl::sycl::detail::modeNeedsOldData
constexpr bool modeNeedsOldData(access::mode m)
Definition: access.hpp:109
cl::sycl::info::global_mem_cache_type::read_only
@ read_only
__OPENCL_GLOBAL_AS__
#define __OPENCL_GLOBAL_AS__
Definition: access.hpp:131
cl::sycl::detail::remove_AS::type
T type
Definition: access.hpp:209
cl::sycl::access::address_space::private_space
@ private_space
cl::sycl::detail::DecoratedType
Definition: access.hpp:162
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12
cl::sycl::detail::isTargetHostAccess
constexpr bool isTargetHostAccess(access::target T)
Definition: access.hpp:105