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