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 __SYCL_EXPORT void waitEvents(std::vector<sycl::event> DepEvents);
46 
47 __SYCL_EXPORT void
48 markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl);
49 
50 template <typename T> T *declptr() { return static_cast<T *>(nullptr); }
51 
52 // Function to get or store id, item, nd_item, group for the host implementation
53 // Pass nullptr to get stored object. Pass valid address to store object
54 template <typename T> T get_or_store(const T *obj) {
55  static thread_local auto stored = *obj;
56  if (obj != nullptr) {
57  stored = *obj;
58  }
59  return stored;
60 }
61 
62 class Builder {
63 public:
64  Builder() = delete;
65 
66  template <int Dims>
67  static group<Dims>
68  createGroup(const range<Dims> &Global, const range<Dims> &Local,
69  const range<Dims> &Group, const id<Dims> &Index) {
70  return group<Dims>(Global, Local, Group, Index);
71  }
72 
73  template <int Dims>
74  static group<Dims> createGroup(const range<Dims> &Global,
75  const range<Dims> &Local,
76  const id<Dims> &Index) {
77  return group<Dims>(Global, Local, Global / Local, Index);
78  }
79 
80  template <class ResType, typename BitsType>
81  static ResType createSubGroupMask(BitsType Bits, size_t BitsNum) {
82  return ResType(Bits, BitsNum);
83  }
84 
85  template <int Dims, bool WithOffset>
86  static std::enable_if_t<WithOffset, item<Dims, WithOffset>>
87  createItem(const range<Dims> &Extent, const id<Dims> &Index,
88  const id<Dims> &Offset) {
89  return item<Dims, WithOffset>(Extent, Index, Offset);
90  }
91 
92  template <int Dims, bool WithOffset>
93  static std::enable_if_t<!WithOffset, item<Dims, WithOffset>>
94  createItem(const range<Dims> &Extent, const id<Dims> &Index) {
95  return item<Dims, WithOffset>(Extent, Index);
96  }
97 
98  template <int Dims>
100  const item<Dims, false> &Local,
101  const group<Dims> &Group) {
102  return nd_item<Dims>(Global, Local, Group);
103  }
104 
105  template <int Dims>
107  const item<Dims, false> &Local) {
108  return h_item<Dims>(Global, Local);
109  }
110 
111  template <int Dims>
113  const item<Dims, false> &Local,
114  const range<Dims> &Flex) {
115  return h_item<Dims>(Global, Local, Flex);
116  }
117 
118  template <int Dims, bool WithOffset>
120  const id<Dims> &NextIndex) {
121  Item.MImpl.MIndex = NextIndex;
122  }
123 
124 #ifdef __SYCL_DEVICE_ONLY__
125 
126  template <int N>
127  using is_valid_dimensions = std::integral_constant<bool, (N > 0) && (N < 4)>;
128 
129  template <int Dims> static const id<Dims> getElement(id<Dims> *) {
130  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
131  return __spirv::initGlobalInvocationId<Dims, id<Dims>>();
132  }
133 
134  template <int Dims> static const group<Dims> getElement(group<Dims> *) {
135  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
136  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
137  range<Dims> LocalSize{__spirv::initWorkgroupSize<Dims, range<Dims>>()};
138  range<Dims> GroupRange{__spirv::initNumWorkgroups<Dims, range<Dims>>()};
139  id<Dims> GroupId{__spirv::initWorkgroupId<Dims, id<Dims>>()};
140  return createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
141  }
142 
143  template <int Dims, bool WithOffset>
144  static std::enable_if_t<WithOffset, const item<Dims, WithOffset>> getItem() {
145  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
146  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
147  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
148  id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
149  return createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
150  }
151 
152  template <int Dims, bool WithOffset>
153  static std::enable_if_t<!WithOffset, const item<Dims, WithOffset>> getItem() {
154  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
155  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
156  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
157  return createItem<Dims, false>(GlobalSize, GlobalId);
158  }
159 
160  template <int Dims> static const nd_item<Dims> getElement(nd_item<Dims> *) {
161  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
162  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
163  range<Dims> LocalSize{__spirv::initWorkgroupSize<Dims, range<Dims>>()};
164  range<Dims> GroupRange{__spirv::initNumWorkgroups<Dims, range<Dims>>()};
165  id<Dims> GroupId{__spirv::initWorkgroupId<Dims, id<Dims>>()};
166  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
167  id<Dims> LocalId{__spirv::initLocalInvocationId<Dims, id<Dims>>()};
168  id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
169  group<Dims> Group =
170  createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
171  item<Dims, true> GlobalItem =
172  createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
173  item<Dims, false> LocalItem = createItem<Dims, false>(LocalSize, LocalId);
174  return createNDItem<Dims>(GlobalItem, LocalItem, Group);
175  }
176 
177  template <int Dims, bool WithOffset>
178  static auto getElement(item<Dims, WithOffset> *)
179  -> decltype(getItem<Dims, WithOffset>()) {
180  return getItem<Dims, WithOffset>();
181  }
182 
183  template <int Dims>
184  static auto getNDItem() -> decltype(getElement(declptr<nd_item<Dims>>())) {
185  return getElement(declptr<nd_item<Dims>>());
186  }
187 
188 #endif // __SYCL_DEVICE_ONLY__
189 };
190 
191 inline constexpr __spv::MemorySemanticsMask::Flag
194 }
195 
196 inline constexpr uint32_t
198  const __spv::MemorySemanticsMask LocalScopeMask =
200  // Huge ternary operator below is a workaround for constexpr function
201  // requirement that such function can only contain return statement and
202  // nothing more
203  //
204  // It is equivalent to the following code:
205  //
206  // uint32_t Flags =
207  // static_cast<uint32_t>(__spv::MemorySemanticsMask::SequentiallyConsistent);
208  // switch (AccessSpace) {
209  // case access::fence_space::global_space:
210  // Flags |=
211  // static_cast<uint32_t>(__spv::MemorySemanticsMask::CrossWorkgroupMemory);
212  // break;
213  // case access::fence_space::local_space:
214  // Flags |= static_cast<uint32_t>(LocalScopeMask);
215  // break;
216  // case access::fence_space::global_and_local:
217  // default:
218  // Flags |= static_cast<uint32_t>(
219  // __spv::MemorySemanticsMask::CrossWorkgroupMemory) |
220  // static_cast<uint32_t>(LocalScopeMask);
221  // break;
222  // }
223  // return Flags;
224 
225  return (AccessSpace == access::fence_space::global_space)
226  ? static_cast<uint32_t>(
229  : (AccessSpace == access::fence_space::local_space)
230  ? static_cast<uint32_t>(
232  LocalScopeMask)
233  : /* default: (AccessSpace ==
234  access::fence_space::global_and_local) */
235  static_cast<uint32_t>(
238  LocalScopeMask);
239 }
240 
241 // To ensure loop unrolling is done when processing dimensions.
242 template <size_t... Inds, class F>
243 void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
244  (f(std::integral_constant<size_t, Inds>{}), ...);
245 }
246 
247 template <size_t count, class F> void loop(F &&f) {
248  loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
249 }
250 inline constexpr bool is_power_of_two(int x) { return (x & (x - 1)) == 0; }
251 } // namespace detail
252 
253 } // namespace _V1
254 } // 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:87
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const range< Dims > &Group, const id< Dims > &Index)
Definition: helpers.hpp:68
static ResType createSubGroupMask(BitsType Bits, size_t BitsNum)
Definition: helpers.hpp:81
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local, const range< Dims > &Flex)
Definition: helpers.hpp:112
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const id< Dims > &Index)
Definition: helpers.hpp:74
static void updateItemIndex(sycl::item< Dims, WithOffset > &Item, const id< Dims > &NextIndex)
Definition: helpers.hpp:119
static std::enable_if_t<!WithOffset, item< Dims, WithOffset > > createItem(const range< Dims > &Extent, const id< Dims > &Index)
Definition: helpers.hpp:94
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local)
Definition: helpers.hpp:106
static nd_item< Dims > createNDItem(const item< Dims, true > &Global, const item< Dims, false > &Local, const group< Dims > &Group)
Definition: helpers.hpp:99
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:49
Identifies an instance of the function object executing at each point in an nd_range.
Definition: nd_item.hpp:48
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: range.hpp:26
constexpr bool is_power_of_two(int x)
Definition: helpers.hpp:250
T get_or_store(const T *obj)
Definition: helpers.hpp:54
void markBufferAsInternal(const std::shared_ptr< buffer_impl > &BufImpl)
Definition: helpers.cpp:30
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
Definition: helpers.hpp:192
void loop_impl(std::integer_sequence< size_t, Inds... >, F &&f)
Definition: helpers.hpp:243
void waitEvents(std::vector< sycl::event > DepEvents)
Definition: helpers.cpp:24
void loop(F &&f)
Definition: helpers.hpp:247
autodecltype(x) x
Definition: access.hpp:18
C++ wrapper of extern "C" PI interfaces.