DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
group.hpp
Go to the documentation of this file.
1 //==-------------- group.hpp --- SYCL work group ---------------------------==//
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_ops.hpp>
17 #include <CL/sycl/detail/spirv.hpp>
18 #include <CL/sycl/device_event.hpp>
19 #include <CL/sycl/h_item.hpp>
20 #include <CL/sycl/id.hpp>
21 #include <CL/sycl/memory_enums.hpp>
22 #include <CL/sycl/pointers.hpp>
23 #include <CL/sycl/range.hpp>
24 #include <stdexcept>
25 #include <type_traits>
26 
28 namespace sycl {
29 namespace detail {
30 class Builder;
31 
32 // Implements a barrier accross work items within a work group.
33 static inline void workGroupBarrier() {
34 #ifdef __SYCL_DEVICE_ONLY__
35  constexpr uint32_t flags =
36  static_cast<uint32_t>(
38  static_cast<uint32_t>(__spv::MemorySemanticsMask::WorkgroupMemory);
40  flags);
41 #endif // __SYCL_DEVICE_ONLY__
42 }
43 
44 } // namespace detail
45 
46 // SYCL 1.2.1rev5, section "4.8.5.3 Parallel For hierarchical invoke":
47 // Quote:
48 // ... To guarantee use of private per-work-item memory, the private_memory
49 // class can be used to wrap the data. This class very simply constructs
50 // private data for a given group across the entire group.The id of the
51 // current work-item is passed to any access to grab the correct data.
52 template <typename T, int Dimensions = 1> class private_memory {
53 public:
54  // Construct based directly off the number of work-items
56 #ifndef __SYCL_DEVICE_ONLY__
57  // serial host => one instance per work-group - allocate space for each WI
58  // in the group:
59  Val.reset(new T[G.get_local_range().size()]);
60 #endif // __SYCL_DEVICE_ONLY__
61  (void)G;
62  }
63 
64  // Access the instance for the current work-item
66 #ifndef __SYCL_DEVICE_ONLY__
67  // Calculate the linear index of current WI and return reference to the
68  // corresponding spot in the value array:
69  size_t Ind = Id.get_physical_local().get_linear_id();
70  return Val.get()[Ind];
71 #else
72  (void)Id;
73  return Val;
74 #endif // __SYCL_DEVICE_ONLY__
75  }
76 
77 private:
78 #ifdef __SYCL_DEVICE_ONLY__
79  // On SYCL device private_memory<T> instance is created per physical WI, so
80  // there is 1:1 correspondence betwen this class instances and per-WI memory.
81  T Val;
82 #else
83  // On serial host there is one private_memory<T> instance per work group, so
84  // it must have space to hold separate value per WI in the group.
85  std::unique_ptr<T[]> Val;
86 #endif // #ifdef __SYCL_DEVICE_ONLY__
87 };
88 
93 template <int Dimensions = 1> class group {
94 public:
95 #ifndef __DISABLE_SYCL_INTEL_GROUP_ALGORITHMS__
98  using linear_id_type = size_t;
99  static constexpr int dimensions = Dimensions;
100 #endif // __DISABLE_SYCL_INTEL_GROUP_ALGORITHMS__
101 
102  static constexpr sycl::memory_scope fence_scope =
103  sycl::memory_scope::work_group;
104 
105  group() = delete;
106 
107  id<Dimensions> get_id() const { return index; }
108 
109  size_t get_id(int dimension) const { return index[dimension]; }
110 
111  range<Dimensions> get_global_range() const { return globalRange; }
112 
113  size_t get_global_range(int dimension) const {
114  return globalRange[dimension];
115  }
116 
117  range<Dimensions> get_local_range() const { return localRange; }
118 
119  size_t get_local_range(int dimension) const { return localRange[dimension]; }
120 
121  range<Dimensions> get_group_range() const { return groupRange; }
122 
123  size_t get_group_range(int dimension) const {
124  return get_group_range()[dimension];
125  }
126 
127  size_t operator[](int dimension) const { return index[dimension]; }
128 
129  template <int dims = Dimensions>
130  typename detail::enable_if_t<(dims == 1), size_t> get_linear_id() const {
131  return index[0];
132  }
133 
134  template <int dims = Dimensions>
135  typename detail::enable_if_t<(dims == 2), size_t> get_linear_id() const {
136  return index[0] * groupRange[1] + index[1];
137  }
138 
139  // SYCL specification 1.2.1rev5, section 4.7.6.5 "Buffer accessor":
140  // Whenever a multi-dimensional index is passed to a SYCL accessor the
141  // linear index is calculated based on the index {id1, id2, id3} provided
142  // and the range of the SYCL accessor {r1, r2, r3} according to row-major
143  // ordering as follows:
144  // id3 + (id2 · r3) + (id1 · r3 · r2) (4.3)
145  // section 4.8.1.8 "group class":
146  // size_t get_linear_id()const
147  // Get a linearized version of the work-group id. Calculating a linear
148  // work-group id from a multi-dimensional index follows the equation 4.3.
149  template <int dims = Dimensions>
150  typename detail::enable_if_t<(dims == 3), size_t> get_linear_id() const {
151  return (index[0] * groupRange[1] * groupRange[2]) +
152  (index[1] * groupRange[2]) + index[2];
153  }
154 
155  template <typename WorkItemFunctionT>
156  void parallel_for_work_item(WorkItemFunctionT Func) const {
157  // need barriers to enforce SYCL semantics for the work item loop -
158  // compilers are expected to optimize when possible
160 #ifdef __SYCL_DEVICE_ONLY__
161  range<Dimensions> GlobalSize{
162  __spirv::initGlobalSize<Dimensions, range<Dimensions>>()};
163  range<Dimensions> LocalSize{
164  __spirv::initWorkgroupSize<Dimensions, range<Dimensions>>()};
165  id<Dimensions> GlobalId{
166  __spirv::initGlobalInvocationId<Dimensions, id<Dimensions>>()};
167  id<Dimensions> LocalId{
168  __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>()};
169 
170  // no 'iterate' in the device code variant, because
171  // (1) this code is already invoked by each work item as a part of the
172  // enclosing parallel_for_work_group kernel
173  // (2) the range this pfwi iterates over matches work group size exactly
174  item<Dimensions, false> GlobalItem =
175  detail::Builder::createItem<Dimensions, false>(GlobalSize, GlobalId);
176  item<Dimensions, false> LocalItem =
177  detail::Builder::createItem<Dimensions, false>(LocalSize, LocalId);
178  h_item<Dimensions> HItem =
179  detail::Builder::createHItem<Dimensions>(GlobalItem, LocalItem);
180 
181  Func(HItem);
182 #else
183  id<Dimensions> GroupStartID = index * localRange;
184 
185  // ... host variant needs explicit 'iterate' because it is serial
187  localRange, [&](const id<Dimensions> &LocalID) {
188  item<Dimensions, false> GlobalItem =
189  detail::Builder::createItem<Dimensions, false>(
190  globalRange, GroupStartID + LocalID);
191  item<Dimensions, false> LocalItem =
192  detail::Builder::createItem<Dimensions, false>(localRange,
193  LocalID);
194  h_item<Dimensions> HItem =
195  detail::Builder::createHItem<Dimensions>(GlobalItem, LocalItem);
196  Func(HItem);
197  });
198 #endif // __SYCL_DEVICE_ONLY__
199  // Need both barriers here - before and after the parallel_for_work_item
200  // (PFWI). There can be work group scope code after the PFWI which reads
201  // work group local data written within this PFWI. Back Ends are expected to
202  // optimize away unneeded barriers (e.g. two barriers in a row).
204  }
205 
206  template <typename WorkItemFunctionT>
208  WorkItemFunctionT Func) const {
210 #ifdef __SYCL_DEVICE_ONLY__
211  range<Dimensions> GlobalSize{
212  __spirv::initGlobalSize<Dimensions, range<Dimensions>>()};
213  range<Dimensions> LocalSize{
214  __spirv::initWorkgroupSize<Dimensions, range<Dimensions>>()};
215  id<Dimensions> GlobalId{
216  __spirv::initGlobalInvocationId<Dimensions, id<Dimensions>>()};
217  id<Dimensions> LocalId{
218  __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>()};
219 
220  item<Dimensions, false> GlobalItem =
221  detail::Builder::createItem<Dimensions, false>(GlobalSize, GlobalId);
222  item<Dimensions, false> LocalItem =
223  detail::Builder::createItem<Dimensions, false>(LocalSize, LocalId);
224  h_item<Dimensions> HItem = detail::Builder::createHItem<Dimensions>(
225  GlobalItem, LocalItem, flexibleRange);
226 
227  // iterate over flexible range with work group size stride; each item
228  // performs flexibleRange/LocalSize iterations (if the former is divisible
229  // by the latter)
231  LocalId, LocalSize, flexibleRange,
232  [&](const id<Dimensions> &LogicalLocalID) {
233  HItem.setLogicalLocalID(LogicalLocalID);
234  Func(HItem);
235  });
236 #else
237  id<Dimensions> GroupStartID = index * localRange;
238 
240  localRange, [&](const id<Dimensions> &LocalID) {
241  item<Dimensions, false> GlobalItem =
242  detail::Builder::createItem<Dimensions, false>(
243  globalRange, GroupStartID + LocalID);
244  item<Dimensions, false> LocalItem =
245  detail::Builder::createItem<Dimensions, false>(localRange,
246  LocalID);
247  h_item<Dimensions> HItem = detail::Builder::createHItem<Dimensions>(
248  GlobalItem, LocalItem, flexibleRange);
249 
251  LocalID, localRange, flexibleRange,
252  [&](const id<Dimensions> &LogicalLocalID) {
253  HItem.setLogicalLocalID(LogicalLocalID);
254  Func(HItem);
255  });
256  });
257 #endif // __SYCL_DEVICE_ONLY__
259  }
260 
263  template <access::mode accessMode = access::mode::read_write>
264  void mem_fence(
265  typename detail::enable_if_t<accessMode == access::mode::read ||
266  accessMode == access::mode::write ||
267  accessMode == access::mode::read_write,
269  accessSpace = access::fence_space::global_and_local) const {
270  uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace);
271  // TODO: currently, there is no good way in SPIR-V to set the memory
272  // barrier only for load operations or only for store operations.
273  // The full read-and-write barrier is used and the template parameter
274  // 'accessMode' is ignored for now. Either SPIR-V or SYCL spec may be
275  // changed to address this discrepancy between SPIR-V and SYCL,
276  // or if we decide that 'accessMode' is the important feature then
277  // we can fix this later, for example, by using OpenCL 1.2 functions
278  // read_mem_fence() and write_mem_fence().
280  }
281 
287  template <typename dataT>
290  size_t numElements, size_t srcStride) const {
291  using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
292  using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
293 
295  __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
296  numElements, srcStride, 0);
297  return device_event(&E);
298  }
299 
305  template <typename dataT>
308  size_t numElements, size_t destStride) const {
309  using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
310  using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
311 
313  __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
314  numElements, destStride, 0);
315  return device_event(&E);
316  }
317 
323  template <typename T, access::address_space DestS, access::address_space SrcS>
326  size_t NumElements, size_t Stride) const {
327  static_assert(sizeof(bool) == sizeof(uint8_t),
328  "Async copy to/from bool memory is not supported.");
329  auto DestP =
330  multi_ptr<uint8_t, DestS>(reinterpret_cast<uint8_t *>(Dest.get()));
331  auto SrcP =
332  multi_ptr<uint8_t, SrcS>(reinterpret_cast<uint8_t *>(Src.get()));
333  return async_work_group_copy(DestP, SrcP, NumElements, Stride);
334  }
335 
341  template <typename T, access::address_space DestS, access::address_space SrcS>
344  size_t NumElements, size_t Stride) const {
345  static_assert(sizeof(bool) == sizeof(uint8_t),
346  "Async copy to/from bool memory is not supported.");
348  auto DestP = multi_ptr<VecT, DestS>(reinterpret_cast<VecT *>(Dest.get()));
349  auto SrcP = multi_ptr<VecT, SrcS>(reinterpret_cast<VecT *>(Src.get()));
350  return async_work_group_copy(DestP, SrcP, NumElements, Stride);
351  }
352 
358  template <typename dataT>
360  global_ptr<dataT> src,
361  size_t numElements) const {
362  return async_work_group_copy(dest, src, numElements, 1);
363  }
364 
370  template <typename dataT>
372  local_ptr<dataT> src,
373  size_t numElements) const {
374  return async_work_group_copy(dest, src, numElements, 1);
375  }
376 
377  template <typename... eventTN>
378  void wait_for(eventTN... Events) const {
379  waitForHelper(Events...);
380  }
381 
382  bool operator==(const group<Dimensions> &rhs) const {
383  bool Result = (rhs.globalRange == globalRange) &&
384  (rhs.localRange == localRange) && (rhs.index == index);
385  __SYCL_ASSERT(rhs.groupRange == groupRange &&
386  "inconsistent group class fields");
387  return Result;
388  }
389 
390  bool operator!=(const group<Dimensions> &rhs) const {
391  return !((*this) == rhs);
392  }
393 
394 private:
395  range<Dimensions> globalRange;
396  range<Dimensions> localRange;
397  range<Dimensions> groupRange;
398  id<Dimensions> index;
399 
400  void waitForHelper() const {}
401 
402  void waitForHelper(device_event Event) const {
403  Event.wait();
404  }
405 
406  template <typename T, typename... Ts>
407  void waitForHelper(T E, Ts... Es) const {
408  waitForHelper(E);
409  waitForHelper(Es...);
410  }
411 
412 protected:
413  friend class detail::Builder;
415  const range<Dimensions> GroupRange, const id<Dimensions> &I)
416  : globalRange(G), localRange(L), groupRange(GroupRange), index(I) {
417  // Make sure local range divides global without remainder:
418  __SYCL_ASSERT(((G % L).size() == 0) &&
419  "global range is not multiple of local");
420  __SYCL_ASSERT((((G / L) - GroupRange).size() == 0) &&
421  "inconsistent group constructor arguments");
422  }
423 };
424 
425 namespace detail {
426 template <int Dims> group<Dims> store_group(const group<Dims> *g) {
427  return get_or_store(g);
428 }
429 } // namespace detail
430 
431 template <int Dims>
432 __SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_group() instead")
433 group<Dims> this_group() {
434 #ifdef __SYCL_DEVICE_ONLY__
435  return detail::Builder::getElement(detail::declptr<group<Dims>>());
436 #else
437  return detail::store_group<Dims>(nullptr);
438 #endif
439 }
440 
441 namespace ext {
442 namespace oneapi {
443 namespace experimental {
444 template <int Dims> group<Dims> this_group() {
445 #ifdef __SYCL_DEVICE_ONLY__
446  return sycl::detail::Builder::getElement(detail::declptr<group<Dims>>());
447 #else
448  return sycl::detail::store_group<Dims>(nullptr);
449 #endif
450 }
451 } // namespace experimental
452 } // namespace oneapi
453 } // namespace ext
454 } // namespace sycl
455 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::group::async_work_group_copy
detail::enable_if_t<!detail::is_bool< dataT >::value, device_event > async_work_group_copy(local_ptr< dataT > dest, global_ptr< dataT > src, size_t numElements, size_t srcStride) const
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
Definition: group.hpp:289
spirv_ops.hpp
cl::sycl::group::get_local_range
size_t get_local_range(int dimension) const
Definition: group.hpp:119
__spirv_ControlBarrier
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:26
cl::sycl::group::mem_fence
void mem_fence(typename detail::enable_if_t< accessMode==access::mode::read||accessMode==access::mode::write||accessMode==access::mode::read_write, access::fence_space > accessSpace=access::fence_space::global_and_local) const
Executes a work-group mem-fence with memory ordering on the local address space, global address space...
Definition: group.hpp:264
cl::sycl::detail::Builder
Definition: helpers.hpp:68
cl::sycl::group::get_linear_id
detail::enable_if_t<(dims==1), size_t > get_linear_id() const
Definition: group.hpp:130
__spv::Scope::Workgroup
@ Workgroup
Definition: spirv_types.hpp:28
T
cl::sycl::item
Identifies an instance of the function object executing at each point in a range.
Definition: helpers.hpp:28
cl::sycl::group::async_work_group_copy
detail::enable_if_t<!detail::is_bool< dataT >::value, device_event > async_work_group_copy(global_ptr< dataT > dest, local_ptr< dataT > src, size_t numElements, size_t destStride) const
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
Definition: group.hpp:307
cl::sycl::group::get_id
size_t get_id(int dimension) const
Definition: group.hpp:109
cl::sycl::private_memory::operator()
T & operator()(const h_item< Dimensions > &Id)
Definition: group.hpp:65
cl::sycl::multi_ptr::get
pointer_t get() const
Definition: multi_ptr.hpp:213
cl::sycl::group::operator!=
bool operator!=(const group< Dimensions > &rhs) const
Definition: group.hpp:390
__SYCL_OpGroupAsyncCopyGlobalToLocal
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest, dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:653
cl::sycl::id< Dimensions >
cl::sycl::group::get_group_range
size_t get_group_range(int dimension) const
Definition: group.hpp:123
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:47
cl::sycl::detail::workGroupBarrier
static void workGroupBarrier()
Definition: group.hpp:33
cl::sycl::group
Encapsulates all functionality required to represent a particular work-group within a parallel execut...
Definition: helpers.hpp:29
cl::sycl::group::get_id
id< Dimensions > get_id() const
Definition: group.hpp:107
cl::sycl::detail::declptr
T * declptr()
Definition: helpers.hpp:56
cl::sycl::group::operator[]
size_t operator[](int dimension) const
Definition: group.hpp:127
cl::sycl::h_item::setLogicalLocalID
void setLogicalLocalID(const id< dimensions > &ID)
Definition: h_item.hpp:124
helpers.hpp
spirv_vars.hpp
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:32
__ocl_event_t
void * __ocl_event_t
Definition: spirv_types.hpp:144
cl::sycl::h_item::get_physical_local
item< dimensions, false > get_physical_local() const
Definition: h_item.hpp:43
cl::sycl::group::wait_for
void wait_for(eventTN... Events) const
Definition: group.hpp:378
cl::sycl::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
id.hpp
cl::sycl::range< Dimensions >
cl::sycl::group::async_work_group_copy
device_event async_work_group_copy(global_ptr< dataT > dest, local_ptr< dataT > src, size_t numElements) const
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
Definition: group.hpp:371
cl::sycl::group::async_work_group_copy
detail::enable_if_t< detail::is_vector_bool< T >::value, device_event > async_work_group_copy(multi_ptr< T, DestS > Dest, multi_ptr< T, SrcS > Src, size_t NumElements, size_t Stride) const
Specialization for vector bool type.
Definition: group.hpp:343
cl::sycl::group::operator==
bool operator==(const group< Dimensions > &rhs) const
Definition: group.hpp:382
cl::sycl::private_memory
Definition: group.hpp:52
cl::sycl::access::fence_space
fence_space
Definition: access.hpp:37
cl::sycl::detail::ConvertToOpenCLType_t
conditional_t< TryToGetVectorT< SelectMatchingOpenCLType_t< T > >::value, typename TryToGetVectorT< SelectMatchingOpenCLType_t< T > >::type, conditional_t< TryToGetPointerT< SelectMatchingOpenCLType_t< T > >::value, typename TryToGetPointerVecT< SelectMatchingOpenCLType_t< T > >::type, SelectMatchingOpenCLType_t< T > >> ConvertToOpenCLType_t
Definition: generic_type_traits.hpp:472
cl::sycl::memory_scope
memory_scope
Definition: memory_enums.hpp:24
cl::sycl::detail::NDLoop
Generates an NDIMS-dimensional perfect loop nest.
Definition: common.hpp:266
cl::sycl::detail::getSPIRVMemorySemanticsMask
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
Definition: helpers.hpp:200
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::store_group
group< Dims > store_group(const group< Dims > *g)
Definition: group.hpp:426
generic_type_traits.hpp
range.hpp
spirv.hpp
cl::sycl::group< 1 >::linear_id_type
size_t linear_id_type
Definition: group.hpp:98
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
__spirv_MemoryBarrier
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:36
cl::sycl::group::get_linear_id
detail::enable_if_t<(dims==3), size_t > get_linear_id() const
Definition: group.hpp:150
cl::sycl::group::async_work_group_copy
device_event async_work_group_copy(local_ptr< dataT > dest, global_ptr< dataT > src, size_t numElements) const
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
Definition: group.hpp:359
cl::sycl::group::get_global_range
size_t get_global_range(int dimension) const
Definition: group.hpp:113
cl::sycl::group::group
group(const range< Dimensions > &G, const range< Dimensions > &L, const range< Dimensions > GroupRange, const id< Dimensions > &I)
Definition: group.hpp:414
__SYCL_OpGroupAsyncCopyLocalToGlobal
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag, dataT *Dest, dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:665
cl::sycl::private_memory::private_memory
private_memory(const group< Dimensions > &G)
Definition: group.hpp:55
__spv::MemorySemanticsMask::SequentiallyConsistent
@ SequentiallyConsistent
Definition: spirv_types.hpp:87
cl::sycl::item::get_linear_id
size_t __SYCL_ALWAYS_INLINE get_linear_id() const
Definition: item.hpp:95
cl::sycl::group::get_linear_id
detail::enable_if_t<(dims==2), size_t > get_linear_id() const
Definition: group.hpp:135
cl::sycl::group::get_local_range
range< Dimensions > get_local_range() const
Definition: group.hpp:117
cl::sycl::detail::change_base_type_t
typename change_base_type< T, B >::type change_base_type_t
Definition: type_traits.hpp:110
__spv::MemorySemanticsMask::WorkgroupMemory
@ WorkgroupMemory
Definition: spirv_types.hpp:90
cl::sycl::device_event
Encapsulates a single SYCL device event which is available only within SYCL kernel functions and can ...
Definition: device_event.hpp:22
cl::sycl::detail::get_or_store
T get_or_store(const T *obj)
Definition: helpers.hpp:60
cl::sycl::group::get_group_range
range< Dimensions > get_group_range() const
Definition: group.hpp:121
cl::sycl::group::async_work_group_copy
detail::enable_if_t< detail::is_scalar_bool< T >::value, device_event > async_work_group_copy(multi_ptr< T, DestS > Dest, multi_ptr< T, SrcS > Src, size_t NumElements, size_t Stride) const
Specialization for scalar bool type.
Definition: group.hpp:325
device_event.hpp
cl::sycl::group::get_global_range
range< Dimensions > get_global_range() const
Definition: group.hpp:111
memory_enums.hpp
cl::sycl::group::parallel_for_work_item
void parallel_for_work_item(WorkItemFunctionT Func) const
Definition: group.hpp:156
common.hpp
h_item.hpp
__SYCL_ASSERT
#define __SYCL_ASSERT(x)
Definition: common.hpp:93
cl::sycl::ext::oneapi::experimental::this_group
group< Dims > this_group()
Definition: group.hpp:444
cl::sycl::group::parallel_for_work_item
void parallel_for_work_item(range< Dimensions > flexibleRange, WorkItemFunctionT Func) const
Definition: group.hpp:207
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
pointers.hpp
spirv_types.hpp
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12