DPC++ Runtime
Runtime libraries for oneAPI DPC++
helpers.hpp
Go to the documentation of this file.
1 //==---------------- helpers.hpp - SYCL helpers ----------------------------==//
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 <CL/__spirv/spirv_types.hpp> // for MemorySemanticsMask
12 #include <sycl/access/access.hpp> // for fence_space
13 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
14 #include <sycl/detail/pi.hpp> // for PiEvent
15 #include <sycl/memory_enums.hpp> // for memory_order
16 
17 #ifdef __SYCL_DEVICE_ONLY__
19 #endif
20 
21 #include <cstddef> // for size_t
22 #include <memory> // for shared_ptr
23 #include <stdint.h> // for uint32_t
24 #include <type_traits> // for enable_if_t, integral_constant
25 #include <utility> // for forward, integer_sequence, mak...
26 #include <vector> // for vector
27 
28 namespace sycl {
29 inline namespace _V1 {
30 class context;
31 class event;
32 template <int Dims, bool WithOffset> class item;
33 template <int Dims> class group;
34 template <int Dims> class range;
35 template <int Dims> class id;
36 template <int Dims> class nd_item;
37 template <int Dims> class h_item;
38 template <typename Type, std::size_t NumElements> class marray;
39 enum class memory_order;
40 
41 namespace detail {
42 
43 class buffer_impl;
44 class context_impl;
45 // The function returns list of events that can be passed to OpenCL API as
46 // dependency list and waits for others.
47 __SYCL_EXPORT std::vector<sycl::detail::pi::PiEvent>
48 getOrWaitEvents(std::vector<sycl::event> DepEvents,
49  std::shared_ptr<sycl::detail::context_impl> Context);
50 
51 __SYCL_EXPORT void waitEvents(std::vector<sycl::event> DepEvents);
52 
53 __SYCL_EXPORT void
54 markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl);
55 
56 template <typename T> T *declptr() { return static_cast<T *>(nullptr); }
57 
58 // Function to get of store id, item, nd_item, group for the host implementation
59 // Pass nullptr to get stored object. Pass valid address to store object
60 template <typename T> T get_or_store(const T *obj) {
61  static thread_local auto stored = *obj;
62  if (obj != nullptr) {
63  stored = *obj;
64  }
65  return stored;
66 }
67 
68 class Builder {
69 public:
70  Builder() = delete;
71 
72  template <int Dims>
73  static group<Dims>
74  createGroup(const range<Dims> &Global, const range<Dims> &Local,
75  const range<Dims> &Group, const id<Dims> &Index) {
76  return group<Dims>(Global, Local, Group, Index);
77  }
78 
79  template <int Dims>
80  static group<Dims> createGroup(const range<Dims> &Global,
81  const range<Dims> &Local,
82  const id<Dims> &Index) {
83  return group<Dims>(Global, Local, Global / Local, Index);
84  }
85 
86  template <class ResType, typename BitsType>
87  static ResType createSubGroupMask(BitsType Bits, size_t BitsNum) {
88  return ResType(Bits, BitsNum);
89  }
90 
91  template <int Dims, bool WithOffset>
92  static std::enable_if_t<WithOffset, item<Dims, WithOffset>>
93  createItem(const range<Dims> &Extent, const id<Dims> &Index,
94  const id<Dims> &Offset) {
95  return item<Dims, WithOffset>(Extent, Index, Offset);
96  }
97 
98  template <int Dims, bool WithOffset>
99  static std::enable_if_t<!WithOffset, item<Dims, WithOffset>>
100  createItem(const range<Dims> &Extent, const id<Dims> &Index) {
101  return item<Dims, WithOffset>(Extent, Index);
102  }
103 
104  template <int Dims>
106  const item<Dims, false> &Local,
107  const group<Dims> &Group) {
108  return nd_item<Dims>(Global, Local, Group);
109  }
110 
111  template <int Dims>
113  const item<Dims, false> &Local) {
114  return h_item<Dims>(Global, Local);
115  }
116 
117  template <int Dims>
119  const item<Dims, false> &Local,
120  const range<Dims> &Flex) {
121  return h_item<Dims>(Global, Local, Flex);
122  }
123 
124  template <int Dims, bool WithOffset>
126  const id<Dims> &NextIndex) {
127  Item.MImpl.MIndex = NextIndex;
128  }
129 
130 #ifdef __SYCL_DEVICE_ONLY__
131 
132  template <int N>
133  using is_valid_dimensions = std::integral_constant<bool, (N > 0) && (N < 4)>;
134 
135  template <int Dims> static const id<Dims> getElement(id<Dims> *) {
136  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
137  return __spirv::initGlobalInvocationId<Dims, id<Dims>>();
138  }
139 
140  template <int Dims> static const group<Dims> getElement(group<Dims> *) {
141  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
142  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
143  range<Dims> LocalSize{__spirv::initWorkgroupSize<Dims, range<Dims>>()};
144  range<Dims> GroupRange{__spirv::initNumWorkgroups<Dims, range<Dims>>()};
145  id<Dims> GroupId{__spirv::initWorkgroupId<Dims, id<Dims>>()};
146  return createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
147  }
148 
149  template <int Dims, bool WithOffset>
150  static std::enable_if_t<WithOffset, const item<Dims, WithOffset>> getItem() {
151  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
152  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
153  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
154  id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
155  return createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
156  }
157 
158  template <int Dims, bool WithOffset>
159  static std::enable_if_t<!WithOffset, const item<Dims, WithOffset>> getItem() {
160  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
161  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
162  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
163  return createItem<Dims, false>(GlobalSize, GlobalId);
164  }
165 
166  template <int Dims> static const nd_item<Dims> getElement(nd_item<Dims> *) {
167  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
168  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
169  range<Dims> LocalSize{__spirv::initWorkgroupSize<Dims, range<Dims>>()};
170  range<Dims> GroupRange{__spirv::initNumWorkgroups<Dims, range<Dims>>()};
171  id<Dims> GroupId{__spirv::initWorkgroupId<Dims, id<Dims>>()};
172  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
173  id<Dims> LocalId{__spirv::initLocalInvocationId<Dims, id<Dims>>()};
174  id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
175  group<Dims> Group =
176  createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
177  item<Dims, true> GlobalItem =
178  createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
179  item<Dims, false> LocalItem = createItem<Dims, false>(LocalSize, LocalId);
180  return createNDItem<Dims>(GlobalItem, LocalItem, Group);
181  }
182 
183  template <int Dims, bool WithOffset>
184  static auto getElement(item<Dims, WithOffset> *)
185  -> decltype(getItem<Dims, WithOffset>()) {
186  return getItem<Dims, WithOffset>();
187  }
188 
189  template <int Dims>
190  static auto getNDItem() -> decltype(getElement(declptr<nd_item<Dims>>())) {
191  return getElement(declptr<nd_item<Dims>>());
192  }
193 
194 #endif // __SYCL_DEVICE_ONLY__
195 };
196 
197 inline constexpr __spv::MemorySemanticsMask::Flag
200 }
201 
202 inline constexpr uint32_t
204  const __spv::MemorySemanticsMask LocalScopeMask =
206  // Huge ternary operator below is a workaround for constexpr function
207  // requirement that such function can only contain return statement and
208  // nothing more
209  //
210  // It is equivalent to the following code:
211  //
212  // uint32_t Flags =
213  // static_cast<uint32_t>(__spv::MemorySemanticsMask::SequentiallyConsistent);
214  // switch (AccessSpace) {
215  // case access::fence_space::global_space:
216  // Flags |=
217  // static_cast<uint32_t>(__spv::MemorySemanticsMask::CrossWorkgroupMemory);
218  // break;
219  // case access::fence_space::local_space:
220  // Flags |= static_cast<uint32_t>(LocalScopeMask);
221  // break;
222  // case access::fence_space::global_and_local:
223  // default:
224  // Flags |= static_cast<uint32_t>(
225  // __spv::MemorySemanticsMask::CrossWorkgroupMemory) |
226  // static_cast<uint32_t>(LocalScopeMask);
227  // break;
228  // }
229  // return Flags;
230 
231  return (AccessSpace == access::fence_space::global_space)
232  ? static_cast<uint32_t>(
235  : (AccessSpace == access::fence_space::local_space)
236  ? static_cast<uint32_t>(
238  LocalScopeMask)
239  : /* default: (AccessSpace ==
240  access::fence_space::global_and_local) */
241  static_cast<uint32_t>(
244  LocalScopeMask);
245 }
246 
247 // To ensure loop unrolling is done when processing dimensions.
248 template <size_t... Inds, class F>
249 void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
250  (f(std::integral_constant<size_t, Inds>{}), ...);
251 }
252 
253 template <size_t count, class F> void loop(F &&f) {
254  loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
255 }
256 } // namespace detail
257 
258 } // namespace _V1
259 } // namespace sycl
static std::enable_if_t< WithOffset, item< Dims, WithOffset > > createItem(const range< Dims > &Extent, const id< Dims > &Index, const id< Dims > &Offset)
Definition: helpers.hpp:93
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const range< Dims > &Group, const id< Dims > &Index)
Definition: helpers.hpp:74
static ResType createSubGroupMask(BitsType Bits, size_t BitsNum)
Definition: helpers.hpp:87
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local, const range< Dims > &Flex)
Definition: helpers.hpp:118
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const id< Dims > &Index)
Definition: helpers.hpp:80
static void updateItemIndex(sycl::item< Dims, WithOffset > &Item, const id< Dims > &NextIndex)
Definition: helpers.hpp:125
static std::enable_if_t<!WithOffset, item< Dims, WithOffset > > createItem(const range< Dims > &Extent, const id< Dims > &Index)
Definition: helpers.hpp:100
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local)
Definition: helpers.hpp:112
static nd_item< Dims > createNDItem(const item< Dims, true > &Global, const item< Dims, false > &Local, const group< Dims > &Group)
Definition: helpers.hpp:105
Identifies an instance of a group::parallel_for_work_item function object executing at each point in ...
Definition: h_item.hpp:31
A unique identifier of an item in an index space.
Definition: id.hpp:36
Identifies an instance of the function object executing at each point in a range.
Definition: item.hpp:37
Provides a cross-platform math array class template that works on SYCL devices as well as in host C++...
Definition: marray.hpp:48
Identifies an instance of the function object executing at each point in an nd_range.
Definition: nd_item.hpp:544
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: range.hpp:26
std::vector< sycl::detail::pi::PiEvent > getOrWaitEvents(std::vector< sycl::event > DepEvents, std::shared_ptr< sycl::detail::context_impl > Context)
T get_or_store(const T *obj)
Definition: helpers.hpp:60
void markBufferAsInternal(const std::shared_ptr< buffer_impl > &BufImpl)
Definition: helpers.cpp:69
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
Definition: helpers.hpp:198
void loop_impl(std::integer_sequence< size_t, Inds... >, F &&f)
Definition: helpers.hpp:249
void waitEvents(std::vector< sycl::event > DepEvents)
Definition: helpers.cpp:63
void loop(F &&f)
Definition: helpers.hpp:253
Definition: access.hpp:18
C++ wrapper of extern "C" PI interfaces.