DPC++ Runtime
Runtime libraries for oneAPI DPC++
device_global.hpp
Go to the documentation of this file.
1 //==----------- device_global.hpp - SYCL device_global extension -----------==//
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 
9 #pragma once
10 
11 #include <sycl/access/access.hpp> // for address_space
12 #include <sycl/exception.hpp> // for make_error_code
13 #include <sycl/ext/oneapi/device_global/properties.hpp> // for device_image...
14 #include <sycl/ext/oneapi/properties/properties.hpp> // for properties_t
15 #include <sycl/multi_ptr.hpp> // for multi_ptr
16 #include <sycl/pointers.hpp> // for decorated_gl...
17 
18 #include <cstddef> // for ptrdiff_t
19 #include <type_traits> // for enable_if_t
20 #include <utility> // for declval
21 
22 #ifdef __SYCL_DEVICE_ONLY__
23 #define __SYCL_HOST_NOT_SUPPORTED(Op)
24 #else
25 #define __SYCL_HOST_NOT_SUPPORTED(Op) \
26  throw sycl::exception( \
27  sycl::make_error_code(sycl::errc::feature_not_supported), \
28  Op " is not supported on host device.");
29 #endif
30 
31 // Helper macro for conditional device_global property meta info filtering. This
32 // lets us ignore certain properties under specified conditions, e.g. ignoring
33 // host_access if device_image_scope isn't also present.
34 #define __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props) \
35  detail::ConditionalPropertyMetaInfo< \
36  Props, detail::DeviceGlobalMetaInfoFilter< \
37  Props, detail::properties_t<Props...>>::value>
38 
39 namespace sycl {
40 inline namespace _V1 {
41 namespace ext::oneapi::experimental {
42 
43 namespace detail {
44 // Type-trait for checking if a type defines `operator->`.
45 template <typename T, typename = void>
46 struct HasArrowOperator : std::false_type {};
47 template <typename T>
49  std::void_t<decltype(std::declval<T>().operator->())>>
50  : std::true_type {};
51 
52 // Base class for device_global.
53 template <typename T, typename PropertyListT, typename = void>
55 protected:
57 
58  // The pointer member is mutable to avoid the compiler optimizing it out when
59  // accessing const-qualified device_global variables.
60  mutable pointer_t usmptr{};
61  const T init_val{};
62 
64  pointer_t get_ptr() const noexcept { return usmptr; }
65 
66 public:
67 #if __cpp_consteval
68  template <typename... Args>
69  consteval explicit device_global_base(Args &&...args) : init_val{args...} {}
70 #else
71  device_global_base() = default;
72 #endif // __cpp_consteval
73 
74  template <access::decorated IsDecorated>
77  __SYCL_HOST_NOT_SUPPORTED("get_multi_ptr()")
79  get_ptr()};
80  }
81 
82  template <access::decorated IsDecorated>
85  __SYCL_HOST_NOT_SUPPORTED("get_multi_ptr()")
87  get_ptr()};
88  }
89 };
90 
91 // Specialization of device_global base class for when device_image_scope is in
92 // the property list.
93 template <typename T, typename... Props>
95  T, properties_t<Props...>,
96  std::enable_if_t<properties_t<Props...>::template has_property<
97  device_image_scope_key>()>> {
98 protected:
99  T val{};
100  T *get_ptr() noexcept { return &val; }
101  const T *get_ptr() const noexcept { return &val; }
102 
103 public:
104 #if __cpp_consteval
105  template <typename... Args>
106  consteval explicit device_global_base(Args &&...args) : val{args...} {}
107 #else
108  device_global_base() = default;
109 #endif // __cpp_consteval
110 
111  template <access::decorated IsDecorated>
114  __SYCL_HOST_NOT_SUPPORTED("get_multi_ptr()")
115  return address_space_cast<access::address_space::global_space, IsDecorated,
116  T>(this->get_ptr());
117  }
118 
119  template <access::decorated IsDecorated>
122  __SYCL_HOST_NOT_SUPPORTED("get_multi_ptr()")
123  return address_space_cast<access::address_space::global_space, IsDecorated,
124  const T>(this->get_ptr());
125  }
126 };
127 } // namespace detail
128 
129 template <typename T, typename PropertyListT = empty_properties_t>
130 class
131 #ifdef __SYCL_DEVICE_ONLY__
132  // FIXME: Temporary work-around. Remove when fixed.
133  [[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global]]
134 #endif
135  device_global {
136  // This should always fail when instantiating the unspecialized version.
138  "Property list is invalid.");
139 };
140 
141 template <typename T, typename... Props>
142 class
143 #ifdef __SYCL_DEVICE_ONLY__
144  [[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global,
145  __sycl_detail__::add_ir_attributes_global_variable(
146  "sycl-device-global-size",
147  __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::name..., sizeof(T),
148  __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::value...)]]
149 #endif
150  device_global<T, detail::properties_t<Props...>>
151  : public detail::device_global_base<T, detail::properties_t<Props...>> {
152 
153  using property_list_t = detail::properties_t<Props...>;
154 
155 public:
156  using element_type = std::remove_extent_t<T>;
157 
158 #if !__cpp_consteval
159  static_assert(std::is_trivially_default_constructible_v<T>,
160  "Type T must be trivially default constructable (until C++20 "
161  "consteval is supported and enabled.)");
162 #endif // !__cpp_consteval
163  static_assert(std::is_trivially_destructible_v<T>,
164  "Type T must be trivially destructible.");
165 
167  "Property list is invalid.");
168 
169  // Inherit the base class' constructors
171  T, detail::properties_t<Props...>>::device_global_base;
172 
173  device_global(const device_global &) = delete;
174  device_global(const device_global &&) = delete;
176  device_global &operator=(const device_global &&) = delete;
177 
178  T &get() noexcept {
180  return *this->get_ptr();
181  }
182 
183  const T &get() const noexcept {
185  return *this->get_ptr();
186  }
187 
188  operator T &() noexcept {
189  __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of device_global to T")
190  return get();
191  }
192 
193  operator const T &() const noexcept {
194  __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of device_global to T")
195  return get();
196  }
197 
198  device_global &operator=(const T &newValue) noexcept {
199  __SYCL_HOST_NOT_SUPPORTED("Assignment operator")
200  *this->get_ptr() = newValue;
201  return *this;
202  }
203 
204  template <class RelayT = T>
205  std::remove_reference_t<
206  decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> &
207  operator[](std::ptrdiff_t idx) noexcept {
208  __SYCL_HOST_NOT_SUPPORTED("Subscript operator")
209  return (*this->get_ptr())[idx];
210  }
211 
212  template <class RelayT = T>
213  const std::remove_reference_t<
214  decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> &
215  operator[](std::ptrdiff_t idx) const noexcept {
216  __SYCL_HOST_NOT_SUPPORTED("Subscript operator")
217  return (*this->get_ptr())[idx];
218  }
219 
220  template <class RelayT = T>
221  std::enable_if_t<detail::HasArrowOperator<RelayT>::value ||
222  std::is_pointer_v<RelayT>,
223  RelayT> &
225  __SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global")
226  return *this->get_ptr();
227  }
228 
229  template <class RelayT = T>
230  std::enable_if_t<detail::HasArrowOperator<RelayT>::value ||
231  std::is_pointer_v<RelayT>,
232  const RelayT> &
234  __SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global")
235  return *this->get_ptr();
236  }
237 
238  template <typename propertyT> static constexpr bool has_property() {
239  return property_list_t::template has_property<propertyT>();
240  }
241 
242  template <typename propertyT> static constexpr auto get_property() {
243  return property_list_t::template get_property<propertyT>();
244  }
245 };
246 
247 } // namespace ext::oneapi::experimental
248 } // namespace _V1
249 } // namespace sycl
250 
251 #undef __SYCL_HOST_NOT_SUPPORTED
252 #undef __SYCL_DEVICE_GLOBAL_PROP_META_INFO
multi_ptr< const T, access::address_space::global_space, IsDecorated > get_multi_ptr() const noexcept
multi_ptr< T, access::address_space::global_space, IsDecorated > get_multi_ptr() noexcept
typename decorated_global_ptr< T >::pointer pointer_t
std::enable_if_t< detail::HasArrowOperator< RelayT >::value||std::is_pointer_v< RelayT >, const RelayT > & operator->() const noexcept
std::remove_reference_t< decltype(std::declval< RelayT >)[std::declval< std::ptrdiff_t >)])> & operator[](std::ptrdiff_t idx) noexcept
std::enable_if_t< detail::HasArrowOperator< RelayT >::value||std::is_pointer_v< RelayT >, RelayT > & operator->() noexcept
const std::remove_reference_t< decltype(std::declval< RelayT >)[std::declval< std::ptrdiff_t >)])> & operator[](std::ptrdiff_t idx) const noexcept
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: multi_ptr.hpp:83
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
Definition: multi_ptr.hpp:95
#define __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)
#define __SYCL_HOST_NOT_SUPPORTED(Op)
Definition: access.hpp:18
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324