DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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 
16 #include <CL/sycl/detail/pi.hpp>
18 
19 #include <memory>
20 #include <stdexcept>
21 #include <type_traits>
22 #include <vector>
23 
25 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 
47 class context_impl;
48 // The function returns list of events that can be passed to OpenCL API as
49 // dependency list and waits for others.
50 __SYCL_EXPORT std::vector<RT::PiEvent>
51 getOrWaitEvents(std::vector<cl::sycl::event> DepEvents,
52  std::shared_ptr<cl::sycl::detail::context_impl> Context);
53 
54 __SYCL_EXPORT void waitEvents(std::vector<cl::sycl::event> DepEvents);
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>
87  static ResType createSubGroupMask(uint32_t Bits, size_t BitsNum) {
88  return ResType(Bits, BitsNum);
89  }
90 
91  template <int Dims, bool 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>
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 detail::enable_if_t<WithOffset, const item<Dims, WithOffset>>
151  getItem() {
152  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
153  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
154  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
155  id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
156  return createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
157  }
158 
159  template <int Dims, bool WithOffset>
160  static detail::enable_if_t<!WithOffset, const item<Dims, WithOffset>>
161  getItem() {
162  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
163  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
164  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
165  return createItem<Dims, false>(GlobalSize, GlobalId);
166  }
167 
168  template <int Dims> static const nd_item<Dims> getElement(nd_item<Dims> *) {
169  static_assert(is_valid_dimensions<Dims>::value, "invalid dimensions");
170  range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
171  range<Dims> LocalSize{__spirv::initWorkgroupSize<Dims, range<Dims>>()};
172  range<Dims> GroupRange{__spirv::initNumWorkgroups<Dims, range<Dims>>()};
173  id<Dims> GroupId{__spirv::initWorkgroupId<Dims, id<Dims>>()};
174  id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
175  id<Dims> LocalId{__spirv::initLocalInvocationId<Dims, id<Dims>>()};
176  id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
177  group<Dims> Group =
178  createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
179  item<Dims, true> GlobalItem =
180  createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
181  item<Dims, false> LocalItem = createItem<Dims, false>(LocalSize, LocalId);
182  return createNDItem<Dims>(GlobalItem, LocalItem, Group);
183  }
184 
185  template <int Dims, bool WithOffset>
186  static auto getElement(item<Dims, WithOffset> *)
187  -> decltype(getItem<Dims, WithOffset>()) {
188  return getItem<Dims, WithOffset>();
189  }
190 
191  template <int Dims>
192  static auto getNDItem() -> decltype(getElement(declptr<nd_item<Dims>>())) {
193  return getElement(declptr<nd_item<Dims>>());
194  }
195 
196 #endif // __SYCL_DEVICE_ONLY__
197 };
198 
199 inline constexpr __spv::MemorySemanticsMask::Flag
202 }
203 
204 inline constexpr uint32_t
206  const __spv::MemorySemanticsMask LocalScopeMask =
208  // Huge ternary operator below is a workaround for constexpr function
209  // requirement that such function can only contain return statement and
210  // nothing more
211  //
212  // It is equivalent to the following code:
213  //
214  // uint32_t Flags =
215  // static_cast<uint32_t>(__spv::MemorySemanticsMask::SequentiallyConsistent);
216  // switch (AccessSpace) {
217  // case access::fence_space::global_space:
218  // Flags |=
219  // static_cast<uint32_t>(__spv::MemorySemanticsMask::CrossWorkgroupMemory);
220  // break;
221  // case access::fence_space::local_space:
222  // Flags |= static_cast<uint32_t>(LocalScopeMask);
223  // break;
224  // case access::fence_space::global_and_local:
225  // default:
226  // Flags |= static_cast<uint32_t>(
227  // __spv::MemorySemanticsMask::CrossWorkgroupMemory) |
228  // static_cast<uint32_t>(LocalScopeMask);
229  // break;
230  // }
231  // return Flags;
232 
233  return (AccessSpace == access::fence_space::global_space)
234  ? static_cast<uint32_t>(
237  : (AccessSpace == access::fence_space::local_space)
238  ? static_cast<uint32_t>(
240  LocalScopeMask)
241  : /* default: (AccessSpace ==
242  access::fence_space::global_and_local) */
243  static_cast<uint32_t>(
246  LocalScopeMask);
247 }
248 
249 } // namespace detail
250 
251 } // namespace sycl
252 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::Builder
Definition: helpers.hpp:68
cl::sycl::memory_order
memory_order
Definition: memory_enums.hpp:14
T
cl::sycl::detail::waitEvents
void waitEvents(std::vector< cl::sycl::event > DepEvents)
Definition: helpers.cpp:36
type_traits.hpp
cl::sycl::detail::Builder::createNDItem
static nd_item< Dims > createNDItem(const item< Dims, true > &Global, const item< Dims, false > &Local, const group< Dims > &Group)
Definition: helpers.hpp:105
cl::sycl::item
Identifies an instance of the function object executing at each point in a range.
Definition: helpers.hpp:28
cl::sycl::detail::getSPIRVMemorySemanticsMask
constexpr uint32_t getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace, const __spv::MemorySemanticsMask LocalScopeMask=__spv::MemorySemanticsMask::WorkgroupMemory)
Definition: helpers.hpp:205
cl::sycl::id
A unique identifier of an item in an index space.
Definition: array.hpp:17
cl::sycl::detail::Builder::createGroup
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const id< Dims > &Index)
Definition: helpers.hpp:80
cl::sycl::group
Encapsulates all functionality required to represent a particular work-group within a parallel execut...
Definition: helpers.hpp:29
cl::sycl::detail::declptr
T * declptr()
Definition: helpers.hpp:56
cl::sycl::detail::Builder::createItem
static detail::enable_if_t<!WithOffset, item< Dims, WithOffset > > createItem(const range< Dims > &Extent, const id< Dims > &Index)
Definition: helpers.hpp:100
spirv_vars.hpp
cl::sycl::detail::Builder::createHItem
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local, const range< Dims > &Flex)
Definition: helpers.hpp:118
__spv::MemorySemanticsMask::CrossWorkgroupMemory
@ CrossWorkgroupMemory
Definition: spirv_types.hpp:91
access.hpp
pi.hpp
cl::sycl::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:24
cl::sycl::access::fence_space
fence_space
Definition: access.hpp:37
export.hpp
__spv::MemorySemanticsMask
Definition: spirv_types.hpp:80
cl::sycl::h_item
Identifies an instance of a group::parallel_for_work_item function object executing at each point in ...
Definition: helpers.hpp:33
cl::sycl::detail::Builder::createGroup
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const range< Dims > &Group, const id< Dims > &Index)
Definition: helpers.hpp:74
cl::sycl::detail::Builder::updateItemIndex
static void updateItemIndex(cl::sycl::item< Dims, WithOffset > &Item, const id< Dims > &NextIndex)
Definition: helpers.hpp:125
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::detail::context_impl
Definition: context_impl.hpp:31
cl::sycl::detail::getOrWaitEvents
std::vector< RT::PiEvent > getOrWaitEvents(std::vector< cl::sycl::event > DepEvents, std::shared_ptr< cl::sycl::detail::context_impl > Context)
__spv::MemorySemanticsMask::SequentiallyConsistent
@ SequentiallyConsistent
Definition: spirv_types.hpp:87
__spv::MemorySemanticsMask::None
@ None
Definition: spirv_types.hpp:83
cl::sycl::detail::Builder::createHItem
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local)
Definition: helpers.hpp:112
__spv::MemorySemanticsMask::WorkgroupMemory
@ WorkgroupMemory
Definition: spirv_types.hpp:90
cl::sycl::detail::get_or_store
T get_or_store(const T *obj)
Definition: helpers.hpp:60
cl::sycl::info::event
event
Definition: info_desc.hpp:279
cl::sycl::nd_item
Identifies an instance of the function object executing at each point in an nd_range.
Definition: helpers.hpp:32
cl::sycl::detail::Builder::createSubGroupMask
static ResType createSubGroupMask(uint32_t Bits, size_t BitsNum)
Definition: helpers.hpp:87
cl::sycl::marray
Provides a cross-patform math array class template that works on SYCL devices as well as in host C++ ...
Definition: helpers.hpp:34
__spv::MemorySemanticsMask::Flag
Flag
Definition: spirv_types.hpp:82
common.hpp
cl::sycl::info::context
context
Definition: info_desc.hpp:40
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::detail::Builder::createItem
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:93
cl::sycl::detail::memcpy
void memcpy(void *Dst, const void *Src, size_t Size)
Definition: helpers.hpp:38
spirv_types.hpp
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12