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