DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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_host_device_space' instead") =
58  generic_space = 6, // TODO generic_space address space is not supported yet
59 };
60 
61 } // namespace access
62 
63 using access::target;
65 
66 template <access_mode mode> struct mode_tag_t {
67  explicit mode_tag_t() = default;
68 };
69 
70 template <access_mode mode, target trgt> struct mode_target_tag_t {
71  explicit mode_target_tag_t() = default;
72 };
73 
74 #if __cplusplus > 201402L
75 
76 inline constexpr mode_tag_t<access_mode::read> read_only{};
77 inline constexpr mode_tag_t<access_mode::read_write> read_write{};
78 inline constexpr mode_tag_t<access_mode::write> write_only{};
79 inline constexpr mode_target_tag_t<access_mode::read, target::constant_buffer>
80  read_constant{};
81 
82 #else
83 
84 namespace {
85 
86 constexpr const auto &read_only =
87  sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::read>>::value;
88 constexpr const auto &read_write = sycl::detail::InlineVariableHelper<
89  mode_tag_t<access_mode::read_write>>::value;
90 constexpr const auto &write_only =
91  sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::write>>::value;
92 constexpr const auto &read_constant = sycl::detail::InlineVariableHelper<
93  mode_target_tag_t<access_mode::read, target::constant_buffer>>::value;
94 
95 } // namespace
96 
97 #endif
98 
99 namespace detail {
100 
102  return T == access::target::host_buffer || T == access::target::host_image;
103 }
104 
105 constexpr bool modeNeedsOldData(access::mode m) {
106  return m == access::mode::read || m == access::mode::write ||
107  m == access::mode::read_write || m == access::mode::atomic;
108 }
109 
110 constexpr bool modeWritesNewData(access::mode m) {
111  return m != access::mode::read;
112 }
113 
114 #ifdef __SYCL_DEVICE_ONLY__
115 #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global))
116 #ifdef __ENABLE_USM_ADDR_SPACE__
117 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device))
118 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host))
119 #else
120 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global))
121 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global))
122 #endif // __ENABLE_USM_ADDR_SPACE__
123 #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
124 #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
125 #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
126 #else
127 #define __OPENCL_GLOBAL_AS__
128 #define __OPENCL_GLOBAL_DEVICE_AS__
129 #define __OPENCL_GLOBAL_HOST_AS__
130 #define __OPENCL_LOCAL_AS__
131 #define __OPENCL_CONSTANT_AS__
132 #define __OPENCL_PRIVATE_AS__
133 #endif
134 
135 template <access::target accessTarget> struct TargetToAS {
136  constexpr static access::address_space AS =
137  access::address_space::global_space;
138 };
139 
140 #ifdef __ENABLE_USM_ADDR_SPACE__
141 template <> struct TargetToAS<access::target::device> {
142  constexpr static access::address_space AS =
143  access::address_space::global_device_space;
144 };
145 #endif // __ENABLE_USM_ADDR_SPACE__
146 
147 template <> struct TargetToAS<access::target::local> {
148  constexpr static access::address_space AS =
149  access::address_space::local_space;
150 };
151 
152 template <> struct TargetToAS<access::target::constant_buffer> {
153  constexpr static access::address_space AS =
154  access::address_space::constant_space;
155 };
156 
157 template <typename ElementType, access::address_space addressSpace>
159 
160 template <typename ElementType>
161 struct DecoratedType<ElementType, access::address_space::private_space> {
162  using type = __OPENCL_PRIVATE_AS__ ElementType;
163 };
164 
165 template <typename ElementType>
166 struct DecoratedType<ElementType, access::address_space::generic_space> {
167  using type = ElementType;
168 };
169 
170 template <typename ElementType>
171 struct DecoratedType<ElementType, access::address_space::global_space> {
172  using type = __OPENCL_GLOBAL_AS__ ElementType;
173 };
174 
175 template <typename ElementType>
176 struct DecoratedType<ElementType, access::address_space::global_device_space> {
177  using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType;
178 };
179 
180 template <typename ElementType>
181 struct DecoratedType<ElementType, access::address_space::global_host_space> {
182  using type = __OPENCL_GLOBAL_HOST_AS__ ElementType;
183 };
184 
185 template <typename ElementType>
186 struct DecoratedType<ElementType, access::address_space::constant_space> {
187  // Current implementation of address spaces handling leads to possibility
188  // of emitting incorrect (in terms of OpenCL) address space casts from
189  // constant to generic (and vise-versa). So, global address space is used here
190  // instead of constant to avoid incorrect address space casts in the produced
191  // device code.
192 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
193  using type = const __OPENCL_GLOBAL_AS__ ElementType;
194 #else
195  using type = __OPENCL_GLOBAL_AS__ ElementType;
196 #endif
197 };
198 
199 template <typename ElementType>
200 struct DecoratedType<ElementType, access::address_space::local_space> {
201  using type = __OPENCL_LOCAL_AS__ ElementType;
202 };
203 template <class T> struct remove_AS { typedef T type; };
204 
205 #ifdef __SYCL_DEVICE_ONLY__
206 template <class T> struct deduce_AS {
207  static_assert(!std::is_same<typename detail::remove_AS<T>::type, T>::value,
208  "Only types with address space attributes are supported");
209 };
210 
211 template <class T> struct remove_AS<__OPENCL_GLOBAL_AS__ T> { typedef T type; };
212 
213 #ifdef __ENABLE_USM_ADDR_SPACE__
214 template <class T> struct remove_AS<__OPENCL_GLOBAL_DEVICE_AS__ T> {
215  typedef T type;
216 };
217 
218 template <class T> struct remove_AS<__OPENCL_GLOBAL_HOST_AS__ T> {
219  typedef T type;
220 };
221 
222 template <class T> struct deduce_AS<__OPENCL_GLOBAL_DEVICE_AS__ T> {
223  static const access::address_space value =
224  access::address_space::global_device_space;
225 };
226 
227 template <class T> struct deduce_AS<__OPENCL_GLOBAL_HOST_AS__ T> {
228  static const access::address_space value =
229  access::address_space::global_host_space;
230 };
231 #endif // __ENABLE_USM_ADDR_SPACE__
232 
233 template <class T> struct remove_AS<__OPENCL_PRIVATE_AS__ T> {
234  typedef T type;
235 };
236 
237 template <class T> struct remove_AS<__OPENCL_LOCAL_AS__ T> { typedef T type; };
238 
239 template <class T> struct remove_AS<__OPENCL_CONSTANT_AS__ T> {
240  typedef T type;
241 };
242 
243 template <class T> struct deduce_AS<__OPENCL_GLOBAL_AS__ T> {
244  static const access::address_space value =
245  access::address_space::global_space;
246 };
247 
248 template <class T> struct deduce_AS<__OPENCL_PRIVATE_AS__ T> {
249  static const access::address_space value =
250  access::address_space::private_space;
251 };
252 
253 template <class T> struct deduce_AS<__OPENCL_LOCAL_AS__ T> {
254  static const access::address_space value = access::address_space::local_space;
255 };
256 
257 template <class T> struct deduce_AS<__OPENCL_CONSTANT_AS__ T> {
258  static const access::address_space value =
259  access::address_space::constant_space;
260 };
261 #endif
262 
263 #undef __OPENCL_GLOBAL_AS__
264 #undef __OPENCL_GLOBAL_DEVICE_AS__
265 #undef __OPENCL_GLOBAL_HOST_AS__
266 #undef __OPENCL_LOCAL_AS__
267 #undef __OPENCL_CONSTANT_AS__
268 #undef __OPENCL_PRIVATE_AS__
269 } // namespace detail
270 
271 } // namespace sycl
272 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::DecoratedType< ElementType, access::address_space::generic_space >::type
ElementType type
Definition: access.hpp:167
__OPENCL_GLOBAL_HOST_AS__
#define __OPENCL_GLOBAL_HOST_AS__
Definition: access.hpp:129
cl::sycl::access::address_space::ext_intel_host_device_space
@ ext_intel_host_device_space
T
cl::sycl::info::device
device
Definition: info_desc.hpp:50
__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::ext::intel::experimental::type
type
Definition: fpga_utils.hpp:22
cl::sycl::access::target::local
@ local
cl::sycl::access::target::host_image
@ host_image
cl::sycl::access::target::image_array
@ image_array
detail
Definition: pi_opencl.cpp:86
cl::sycl::access::address_space::constant_space
@ constant_space
__OPENCL_LOCAL_AS__
#define __OPENCL_LOCAL_AS__
Definition: access.hpp:130
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:172
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:131
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:35
cl::sycl::access::target::host_buffer
@ host_buffer
cl::sycl::detail::TargetToAS
Definition: access.hpp:135
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::detail::DecoratedType< ElementType, access::address_space::global_host_space >::type
__OPENCL_GLOBAL_HOST_AS__ ElementType type
Definition: access.hpp:182
__OPENCL_GLOBAL_DEVICE_AS__
#define __OPENCL_GLOBAL_DEVICE_AS__
Definition: access.hpp:128
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:162
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:195
cl::sycl::access::address_space
address_space
Definition: access.hpp:45
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:203
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:201
__OPENCL_PRIVATE_AS__
#define __OPENCL_PRIVATE_AS__
Definition: access.hpp:132
cl::sycl::access::mode::discard_write
@ discard_write
cl::sycl::mode_target_tag_t
Definition: access.hpp:70
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:66
cl::sycl::detail::DecoratedType< ElementType, access::address_space::global_device_space >::type
__OPENCL_GLOBAL_DEVICE_AS__ ElementType type
Definition: access.hpp:177
cl::sycl::detail::modeWritesNewData
constexpr bool modeWritesNewData(access::mode m)
Definition: access.hpp:110
cl::sycl::detail::modeNeedsOldData
constexpr bool modeNeedsOldData(access::mode m)
Definition: access.hpp:105
cl::sycl::info::global_mem_cache_type::read_only
@ read_only
__OPENCL_GLOBAL_AS__
#define __OPENCL_GLOBAL_AS__
Definition: access.hpp:127
cl::sycl::detail::remove_AS::type
T type
Definition: access.hpp:203
cl::sycl::atomic
Definition: atomic.hpp:171
cl::sycl::access::address_space::private_space
@ private_space
cl::sycl::detail::DecoratedType
Definition: access.hpp:158
__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:101