DPC++ Runtime
Runtime libraries for oneAPI DPC++
cg_types.hpp
Go to the documentation of this file.
1 //==---- cg_types.hpp - Auxiliary types required by command group class ----==//
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 <sycl/detail/array.hpp> // for array
12 #include <sycl/detail/common.hpp> // for InitializedVal, NDLoop
13 #include <sycl/detail/helpers.hpp> // for Builder
14 #include <sycl/detail/host_profiling_info.hpp> // for HostProfilingInfo
15 #include <sycl/detail/item_base.hpp> // for id
16 #include <sycl/detail/kernel_desc.hpp> // for kernel_param_kind_t
17 #include <sycl/detail/pi.h> // for PI_ERROR_INVALID_WORK...
18 #include <sycl/exception.hpp> // for nd_range_error
19 #include <sycl/group.hpp> // for group
20 #include <sycl/h_item.hpp> // for h_item
21 #include <sycl/id.hpp> // for id
22 #include <sycl/interop_handle.hpp> // for interop_handle
23 #include <sycl/item.hpp> // for item
24 #include <sycl/kernel_handler.hpp> // for kernel_handler
25 #include <sycl/nd_item.hpp> // for nd_item
26 #include <sycl/nd_range.hpp> // for nd_range
27 #include <sycl/range.hpp> // for range, operator*
28 
29 #include <functional> // for function
30 #include <stddef.h> // for size_t
31 #include <type_traits> // for enable_if_t, false_type
32 #include <utility> // for declval
33 
34 namespace sycl {
35 inline namespace _V1 {
36 namespace detail {
37 
38 // The structure represents kernel argument.
39 class ArgDesc {
40 public:
41  ArgDesc(sycl::detail::kernel_param_kind_t Type, void *Ptr, int Size,
42  int Index)
43  : MType(Type), MPtr(Ptr), MSize(Size), MIndex(Index) {}
44 
46  void *MPtr;
47  int MSize;
48  int MIndex;
49 };
50 
51 // The structure represents NDRange - global, local sizes, global offset and
52 // number of dimensions.
53 class NDRDescT {
54  // The method initializes all sizes for dimensions greater than the passed one
55  // to the default values, so they will not affect execution.
56  void setNDRangeLeftover(int Dims_) {
57  for (int I = Dims_; I < 3; ++I) {
58  GlobalSize[I] = 1;
59  LocalSize[I] = LocalSize[0] ? 1 : 0;
60  GlobalOffset[I] = 0;
61  NumWorkGroups[I] = 0;
62  }
63  }
64 
65 public:
67  : GlobalSize{0, 0, 0}, LocalSize{0, 0, 0}, NumWorkGroups{0, 0, 0} {}
68 
69  template <int Dims_> void set(sycl::range<Dims_> NumWorkItems) {
70  for (int I = 0; I < Dims_; ++I) {
71  GlobalSize[I] = NumWorkItems[I];
72  LocalSize[I] = 0;
73  GlobalOffset[I] = 0;
74  NumWorkGroups[I] = 0;
75  }
76  setNDRangeLeftover(Dims_);
77  Dims = Dims_;
78  }
79 
80  // Initializes this ND range descriptor with given range of work items and
81  // offset.
82  template <int Dims_>
83  void set(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset) {
84  for (int I = 0; I < Dims_; ++I) {
85  GlobalSize[I] = NumWorkItems[I];
86  LocalSize[I] = 0;
87  GlobalOffset[I] = Offset[I];
88  NumWorkGroups[I] = 0;
89  }
90  setNDRangeLeftover(Dims_);
91  Dims = Dims_;
92  }
93 
94  template <int Dims_> void set(sycl::nd_range<Dims_> ExecutionRange) {
95  for (int I = 0; I < Dims_; ++I) {
96  GlobalSize[I] = ExecutionRange.get_global_range()[I];
97  LocalSize[I] = ExecutionRange.get_local_range()[I];
98  GlobalOffset[I] = ExecutionRange.get_offset()[I];
99  NumWorkGroups[I] = 0;
100  }
101  setNDRangeLeftover(Dims_);
102  Dims = Dims_;
103  }
104 
105  void set(int Dims_, sycl::nd_range<3> ExecutionRange) {
106  for (int I = 0; I < Dims_; ++I) {
107  GlobalSize[I] = ExecutionRange.get_global_range()[I];
108  LocalSize[I] = ExecutionRange.get_local_range()[I];
109  GlobalOffset[I] = ExecutionRange.get_offset()[I];
110  NumWorkGroups[I] = 0;
111  }
112  setNDRangeLeftover(Dims_);
113  Dims = Dims_;
114  }
115 
116  template <int Dims_> void setNumWorkGroups(sycl::range<Dims_> N) {
117  for (int I = 0; I < Dims_; ++I) {
118  GlobalSize[I] = 0;
119  // '0' is a mark to adjust before kernel launch when there is enough info:
120  LocalSize[I] = 0;
121  GlobalOffset[I] = 0;
122  NumWorkGroups[I] = N[I];
123  }
124  setNDRangeLeftover(Dims_);
125  Dims = Dims_;
126  }
127 
135  size_t Dims;
136 };
137 
138 template <typename, typename T> struct check_fn_signature {
139  static_assert(std::integral_constant<T, false>::value,
140  "Second template parameter is required to be of function type");
141 };
142 
143 template <typename F, typename RetT, typename... Args>
144 struct check_fn_signature<F, RetT(Args...)> {
145 private:
146  template <typename T>
147  static constexpr auto check(T *) -> typename std::is_same<
148  decltype(std::declval<T>().operator()(std::declval<Args>()...)),
149  RetT>::type;
150 
151  template <typename> static constexpr std::false_type check(...);
152 
153  using type = decltype(check<F>(0));
154 
155 public:
156  static constexpr bool value = type::value;
157 };
158 
159 template <typename F, typename... Args>
160 static constexpr bool check_kernel_lambda_takes_args() {
161  return check_fn_signature<std::remove_reference_t<F>, void(Args...)>::value;
162 }
163 
164 // isKernelLambdaCallableWithKernelHandlerImpl checks if LambdaArgType is void
165 // (e.g., in single_task), and based on that, calls
166 // check_kernel_lambda_takes_args with proper set of arguments. Also this type
167 // trait workarounds compilation error which happens only with msvc.
168 
169 template <
170  typename KernelType, typename LambdaArgType,
171  typename std::enable_if_t<std::is_same_v<LambdaArgType, void>> * = nullptr>
173  return check_kernel_lambda_takes_args<KernelType, kernel_handler>();
174 }
175 
176 template <
177  typename KernelType, typename LambdaArgType,
178  typename std::enable_if_t<!std::is_same_v<LambdaArgType, void>> * = nullptr>
180  return check_kernel_lambda_takes_args<KernelType, LambdaArgType,
181  kernel_handler>();
182 }
183 
184 // Type trait to find out if kernal lambda has kernel_handler argument
185 template <typename KernelType, typename LambdaArgType = void>
187  constexpr static bool value =
188  isKernelLambdaCallableWithKernelHandlerImpl<KernelType, LambdaArgType>();
189 };
190 
191 // Helpers for running kernel lambda on the host device
192 
193 template <typename KernelType>
194 typename std::enable_if_t<KernelLambdaHasKernelHandlerArgT<KernelType>::value>
195 runKernelWithoutArg(KernelType KernelName) {
196  kernel_handler KH;
197  KernelName(KH);
198 }
199 
200 template <typename KernelType>
201 typename std::enable_if_t<!KernelLambdaHasKernelHandlerArgT<KernelType>::value>
202 runKernelWithoutArg(KernelType KernelName) {
203  KernelName();
204 }
205 
206 template <typename ArgType, typename KernelType>
207 typename std::enable_if_t<
209 runKernelWithArg(KernelType KernelName, ArgType Arg) {
210  kernel_handler KH;
211  KernelName(Arg, KH);
212 }
213 
214 template <typename ArgType, typename KernelType>
215 typename std::enable_if_t<
217 runKernelWithArg(KernelType KernelName, ArgType Arg) {
218  KernelName(Arg);
219 }
220 
221 // The pure virtual class aimed to store lambda/functors of any type.
223 public:
224  // The method executes lambda stored using NDRange passed.
225  virtual void call(const NDRDescT &NDRDesc, HostProfilingInfo *HPI) = 0;
226  // Return pointer to the lambda object.
227  // Used to extract captured variables.
228  virtual char *getPtr() = 0;
229  virtual ~HostKernelBase() = default;
230 };
231 
232 class HostTask {
233  std::function<void()> MHostTask;
234  std::function<void(interop_handle)> MInteropTask;
235 
236 public:
237  HostTask() : MHostTask([]() {}) {}
238  HostTask(std::function<void()> &&Func) : MHostTask(Func) {}
239  HostTask(std::function<void(interop_handle)> &&Func) : MInteropTask(Func) {}
240 
241  bool isInteropTask() const { return !!MInteropTask; }
242 
243  void call() { MHostTask(); }
244  void call(interop_handle handle) { MInteropTask(handle); }
245 };
246 
247 // Class which stores specific lambda object.
248 template <class KernelType, class KernelArgType, int Dims>
249 class HostKernel : public HostKernelBase {
250  using IDBuilder = sycl::detail::Builder;
251  KernelType MKernel;
252  // Allowing accessing MKernel from 'ResetHostKernelHelper' method of
253  // 'sycl::handler'
254  friend class sycl::handler;
255 
256 public:
257  HostKernel(KernelType Kernel) : MKernel(Kernel) {}
258  void call(const NDRDescT &NDRDesc, HostProfilingInfo *HPI) override {
259  // adjust ND range for serial host:
260  NDRDescT AdjustedRange = NDRDesc;
261 
262  if (NDRDesc.GlobalSize[0] == 0 && NDRDesc.NumWorkGroups[0] != 0) {
263  // This is a special case - NDRange information is not complete, only the
264  // desired number of work groups is set by the user. Choose work group
265  // size (LocalSize), calculate the missing NDRange characteristics
266  // needed to invoke the kernel and adjust the NDRange descriptor
267  // accordingly. For some devices the work group size selection requires
268  // access to the device's properties, hence such late "adjustment".
269  range<3> WGsize{1, 1, 1}; // no better alternative for serial host?
270  AdjustedRange.set(NDRDesc.Dims,
271  nd_range<3>(NDRDesc.NumWorkGroups * WGsize, WGsize));
272  }
273  // If local size for host is not set explicitly, let's adjust it to 1,
274  // so nd_range_error for zero local size is not thrown.
275  if (AdjustedRange.LocalSize[0] == 0)
276  for (size_t I = 0; I < AdjustedRange.Dims; ++I)
277  AdjustedRange.LocalSize[I] = 1;
278  if (HPI)
279  HPI->start();
280  runOnHost(AdjustedRange);
281  if (HPI)
282  HPI->end();
283  }
284 
285  char *getPtr() override { return reinterpret_cast<char *>(&MKernel); }
286 
287  template <class ArgT = KernelArgType>
288  typename std::enable_if_t<std::is_same_v<ArgT, void>>
289  runOnHost(const NDRDescT &) {
290  runKernelWithoutArg(MKernel);
291  }
292 
293  template <class ArgT = KernelArgType>
294  typename std::enable_if_t<std::is_same_v<ArgT, sycl::id<Dims>>>
295  runOnHost(const NDRDescT &NDRDesc) {
297  sycl::id<Dims> Offset;
298  sycl::range<Dims> Stride(
299  InitializedVal<Dims, range>::template get<1>()); // initialized to 1
300  sycl::range<Dims> UpperBound(
302  for (int I = 0; I < Dims; ++I) {
303  Range[I] = NDRDesc.GlobalSize[I];
304  Offset[I] = NDRDesc.GlobalOffset[I];
305  UpperBound[I] = Range[I] + Offset[I];
306  }
307 
309  /*LowerBound=*/Offset, Stride, UpperBound,
310  [&](const sycl::id<Dims> &ID) {
311  sycl::item<Dims, /*Offset=*/true> Item =
312  IDBuilder::createItem<Dims, true>(Range, ID, Offset);
313 
314  runKernelWithArg<const sycl::id<Dims> &>(MKernel, ID);
315  });
316  }
317 
318  template <class ArgT = KernelArgType>
319  typename std::enable_if_t<std::is_same_v<ArgT, item<Dims, /*Offset=*/false>>>
320  runOnHost(const NDRDescT &NDRDesc) {
321  sycl::id<Dims> ID;
323  for (int I = 0; I < Dims; ++I)
324  Range[I] = NDRDesc.GlobalSize[I];
325 
326  detail::NDLoop<Dims>::iterate(Range, [&](const sycl::id<Dims> ID) {
327  sycl::item<Dims, /*Offset=*/false> Item =
328  IDBuilder::createItem<Dims, false>(Range, ID);
329  sycl::item<Dims, /*Offset=*/true> ItemWithOffset = Item;
330 
331  runKernelWithArg<sycl::item<Dims, /*Offset=*/false>>(MKernel, Item);
332  });
333  }
334 
335  template <class ArgT = KernelArgType>
336  typename std::enable_if_t<std::is_same_v<ArgT, item<Dims, /*Offset=*/true>>>
337  runOnHost(const NDRDescT &NDRDesc) {
339  sycl::id<Dims> Offset;
340  sycl::range<Dims> Stride(
341  InitializedVal<Dims, range>::template get<1>()); // initialized to 1
342  sycl::range<Dims> UpperBound(
344  for (int I = 0; I < Dims; ++I) {
345  Range[I] = NDRDesc.GlobalSize[I];
346  Offset[I] = NDRDesc.GlobalOffset[I];
347  UpperBound[I] = Range[I] + Offset[I];
348  }
349 
351  /*LowerBound=*/Offset, Stride, UpperBound,
352  [&](const sycl::id<Dims> &ID) {
353  sycl::item<Dims, /*Offset=*/true> Item =
354  IDBuilder::createItem<Dims, true>(Range, ID, Offset);
355 
356  runKernelWithArg<sycl::item<Dims, /*Offset=*/true>>(MKernel, Item);
357  });
358  }
359 
360  template <class ArgT = KernelArgType>
361  typename std::enable_if_t<std::is_same_v<ArgT, nd_item<Dims>>>
362  runOnHost(const NDRDescT &NDRDesc) {
364  for (int I = 0; I < Dims; ++I) {
365  if (NDRDesc.LocalSize[I] == 0 ||
366  NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] != 0)
367  throw sycl::nd_range_error("Invalid local size for global size",
368  PI_ERROR_INVALID_WORK_GROUP_SIZE);
369  GroupSize[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
370  }
371 
373  sycl::range<Dims> GlobalSize(
375  sycl::id<Dims> GlobalOffset;
376  for (int I = 0; I < Dims; ++I) {
377  GlobalOffset[I] = NDRDesc.GlobalOffset[I];
378  LocalSize[I] = NDRDesc.LocalSize[I];
379  GlobalSize[I] = NDRDesc.GlobalSize[I];
380  }
381 
382  detail::NDLoop<Dims>::iterate(GroupSize, [&](const id<Dims> &GroupID) {
383  sycl::group<Dims> Group = IDBuilder::createGroup<Dims>(
384  GlobalSize, LocalSize, GroupSize, GroupID);
385 
386  detail::NDLoop<Dims>::iterate(LocalSize, [&](const id<Dims> &LocalID) {
387  id<Dims> GlobalID = GroupID * LocalSize + LocalID + GlobalOffset;
388  const sycl::item<Dims, /*Offset=*/true> GlobalItem =
389  IDBuilder::createItem<Dims, true>(GlobalSize, GlobalID,
390  GlobalOffset);
391  const sycl::item<Dims, /*Offset=*/false> LocalItem =
392  IDBuilder::createItem<Dims, false>(LocalSize, LocalID);
393  const sycl::nd_item<Dims> NDItem =
394  IDBuilder::createNDItem<Dims>(GlobalItem, LocalItem, Group);
395 
396  runKernelWithArg<const sycl::nd_item<Dims>>(MKernel, NDItem);
397  });
398  });
399  }
400 
401  template <typename ArgT = KernelArgType>
402  std::enable_if_t<std::is_same_v<ArgT, sycl::group<Dims>>>
403  runOnHost(const NDRDescT &NDRDesc) {
405 
406  for (int I = 0; I < Dims; ++I) {
407  if (NDRDesc.LocalSize[I] == 0 ||
408  NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] != 0)
409  throw sycl::nd_range_error("Invalid local size for global size",
410  PI_ERROR_INVALID_WORK_GROUP_SIZE);
411  NGroups[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
412  }
413 
415  sycl::range<Dims> GlobalSize(
417  for (int I = 0; I < Dims; ++I) {
418  LocalSize[I] = NDRDesc.LocalSize[I];
419  GlobalSize[I] = NDRDesc.GlobalSize[I];
420  }
421  detail::NDLoop<Dims>::iterate(NGroups, [&](const id<Dims> &GroupID) {
422  sycl::group<Dims> Group =
423  IDBuilder::createGroup<Dims>(GlobalSize, LocalSize, NGroups, GroupID);
424  runKernelWithArg<sycl::group<Dims>>(MKernel, Group);
425  });
426  }
427 
428  ~HostKernel() = default;
429 };
430 
431 } // namespace detail
432 } // namespace _V1
433 } // namespace sycl
sycl::_V1::detail::NDRDescT::NumWorkGroups
sycl::range< 3 > NumWorkGroups
Number of workgroups, used to record the number of workgroups from the simplest form of parallel_for_...
Definition: cg_types.hpp:134
sycl::_V1::detail::HostKernelBase::getPtr
virtual char * getPtr()=0
sycl::_V1::detail::InitializedVal
Definition: common.hpp:258
sycl::_V1::detail::HostKernel::call
void call(const NDRDescT &NDRDesc, HostProfilingInfo *HPI) override
Definition: cg_types.hpp:258
pi.h
sycl::_V1::detail::runKernelWithoutArg
std::enable_if_t< KernelLambdaHasKernelHandlerArgT< KernelType >::value > runKernelWithoutArg(KernelType KernelName)
Definition: cg_types.hpp:195
sycl::_V1::detail::HostProfilingInfo::end
void end()
Measures event's end time.
Definition: event_impl.cpp:381
sycl::_V1::detail::NDLoop::iterate
static __SYCL_ALWAYS_INLINE void iterate(const LoopBoundTy< NDims > &UpperBound, FuncTy f)
Generates ND loop nest with {0,..0} .
Definition: common.hpp:325
sycl::_V1::detail::NDRDescT::GlobalSize
sycl::range< 3 > GlobalSize
Definition: cg_types.hpp:128
sycl::_V1::detail::NDRDescT::set
void set(sycl::range< Dims_ > NumWorkItems, sycl::id< Dims_ > Offset)
Definition: cg_types.hpp:83
item.hpp
sycl::_V1::detail::kernel_param_kind_t
kernel_param_kind_t
Definition: kernel_desc.hpp:38
sycl::_V1::detail::get< 0 >
Definition: tuple.hpp:74
sycl::_V1::detail::isKernelLambdaCallableWithKernelHandlerImpl
constexpr bool isKernelLambdaCallableWithKernelHandlerImpl()
Definition: cg_types.hpp:172
sycl::_V1::detail::HostKernel::runOnHost
std::enable_if_t< std::is_same_v< ArgT, item< Dims, false > > > runOnHost(const NDRDescT &NDRDesc)
Definition: cg_types.hpp:320
sycl::_V1::detail::KernelLambdaHasKernelHandlerArgT
Definition: cg_types.hpp:186
sycl::_V1::detail::HostTask::call
void call()
Definition: cg_types.hpp:243
common.hpp
array.hpp
sycl::_V1::detail::runKernelWithArg
std::enable_if_t< KernelLambdaHasKernelHandlerArgT< KernelType, ArgType >::value > runKernelWithArg(KernelType KernelName, ArgType Arg)
Definition: cg_types.hpp:209
detail
---— Error handling, matching OpenCL plugin semantics.
Definition: common.hpp:44
sycl::_V1::detail::KernelLambdaHasKernelHandlerArgT::value
constexpr static bool value
Definition: cg_types.hpp:187
helpers.hpp
sycl::_V1::detail::HostTask::HostTask
HostTask(std::function< void()> &&Func)
Definition: cg_types.hpp:238
sycl
Definition: access.hpp:18
host_profiling_info.hpp
sycl::_V1::detail::HostKernel
Definition: cg_types.hpp:249
sycl::_V1::id
A unique identifier of an item in an index space.
Definition: array.hpp:20
sycl::_V1::detail::ArgDesc::MType
sycl::detail::kernel_param_kind_t MType
Definition: cg_types.hpp:45
group.hpp
sycl::_V1::nd_range::get_offset
id< Dimensions > get_offset() const
Definition: nd_range.hpp:50
sycl::_V1::detail::NDRDescT::set
void set(int Dims_, sycl::nd_range< 3 > ExecutionRange)
Definition: cg_types.hpp:105
id.hpp
sycl::_V1::detail::HostKernelBase
Definition: cg_types.hpp:222
sycl::_V1::detail::HostKernel::runOnHost
std::enable_if_t< std::is_same_v< ArgT, sycl::id< Dims > > > runOnHost(const NDRDescT &NDRDesc)
Definition: cg_types.hpp:295
interop_handle.hpp
sycl::_V1::detail::HostKernel::HostKernel
HostKernel(KernelType Kernel)
Definition: cg_types.hpp:257
sycl::_V1::detail::check_fn_signature
Definition: cg_types.hpp:138
sycl::_V1::detail::HostProfilingInfo::start
void start()
Measures event's start time.
Definition: event_impl.cpp:379
sycl::_V1::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:51
nd_range.hpp
sycl::_V1::detail::ArgDesc::ArgDesc
ArgDesc(sycl::detail::kernel_param_kind_t Type, void *Ptr, int Size, int Index)
Definition: cg_types.hpp:41
sycl::_V1::nd_item
Identifies an instance of the function object executing at each point in an nd_range.
Definition: helpers.hpp:36
sycl::_V1::detail::NDRDescT::setNumWorkGroups
void setNumWorkGroups(sycl::range< Dims_ > N)
Definition: cg_types.hpp:116
sycl::_V1::nd_range::get_global_range
range< Dimensions > get_global_range() const
Definition: nd_range.hpp:43
sycl::_V1::detail::HostKernel::runOnHost
std::enable_if_t< std::is_same_v< ArgT, nd_item< Dims > > > runOnHost(const NDRDescT &NDRDesc)
Definition: cg_types.hpp:362
sycl::_V1::nd_range::get_local_range
range< Dimensions > get_local_range() const
Definition: nd_range.hpp:45
sycl::_V1::item
Identifies an instance of the function object executing at each point in a range.
Definition: helpers.hpp:32
range.hpp
sycl::_V1::detail::HostTask::call
void call(interop_handle handle)
Definition: cg_types.hpp:244
sycl::_V1::detail::HostKernel::runOnHost
std::enable_if_t< std::is_same_v< ArgT, sycl::group< Dims > > > runOnHost(const NDRDescT &NDRDesc)
Definition: cg_types.hpp:403
sycl::_V1::detail::HostKernel::getPtr
char * getPtr() override
Definition: cg_types.hpp:285
sycl::_V1::detail::HostTask
Definition: cg_types.hpp:232
sycl::_V1::detail::NDRDescT::set
void set(sycl::nd_range< Dims_ > ExecutionRange)
Definition: cg_types.hpp:94
sycl::_V1::handler
Command group handler class.
Definition: handler.hpp:454
sycl::_V1::detail::NDLoop
Generates an NDims-dimensional perfect loop nest.
Definition: common.hpp:319
sycl::_V1::detail::HostTask::HostTask
HostTask(std::function< void(interop_handle)> &&Func)
Definition: cg_types.hpp:239
sycl::_V1::nd_range
Defines the iteration domain of both the work-groups and the overall dispatch.
Definition: uniform.hpp:30
sycl::_V1::detail::ArgDesc::MSize
int MSize
Definition: cg_types.hpp:47
sycl::_V1::detail::HostKernel::~HostKernel
~HostKernel()=default
sycl::_V1::detail::NDRDescT
Definition: cg_types.hpp:53
sycl::_V1::detail::HostKernel::runOnHost
std::enable_if_t< std::is_same_v< ArgT, void > > runOnHost(const NDRDescT &)
Definition: cg_types.hpp:289
sycl::_V1::detail::NDRDescT::NDRDescT
NDRDescT()
Definition: cg_types.hpp:66
sycl::_V1::detail::HostKernelBase::~HostKernelBase
virtual ~HostKernelBase()=default
sycl::_V1::detail::NDRDescT::LocalSize
sycl::range< 3 > LocalSize
Definition: cg_types.hpp:129
kernel_handler.hpp
sycl::_V1::detail::NDRDescT::GlobalOffset
sycl::id< 3 > GlobalOffset
Definition: cg_types.hpp:130
exception.hpp
sycl::_V1::detail::HostKernel::runOnHost
std::enable_if_t< std::is_same_v< ArgT, item< Dims, true > > > runOnHost(const NDRDescT &NDRDesc)
Definition: cg_types.hpp:337
kernel_desc.hpp
item_base.hpp
sycl::_V1::detail::ArgDesc
Definition: cg_types.hpp:39
sycl::_V1::detail::HostTask::HostTask
HostTask()
Definition: cg_types.hpp:237
sycl::_V1::detail::HostKernelBase::call
virtual void call(const NDRDescT &NDRDesc, HostProfilingInfo *HPI)=0
h_item.hpp
sycl::_V1::detail::NDRDescT::set
void set(sycl::range< Dims_ > NumWorkItems)
Definition: cg_types.hpp:69
sycl::_V1::detail::HostProfilingInfo
Profiling info for the host execution.
Definition: host_profiling_info.hpp:20
nd_item.hpp
sycl::_V1::detail::HostTask::isInteropTask
bool isInteropTask() const
Definition: cg_types.hpp:241
sycl::_V1::detail::ArgDesc::MPtr
void * MPtr
Definition: cg_types.hpp:46
sycl::_V1::interop_handle
Definition: interop_handle.hpp:48
sycl::_V1::detail::check_kernel_lambda_takes_args
static constexpr bool check_kernel_lambda_takes_args()
Definition: cg_types.hpp:160
sycl::_V1::detail::get
Definition: tuple.hpp:58
sycl::_V1::group
Definition: helpers.hpp:33
sycl::_V1::detail::NDRDescT::Dims
size_t Dims
Definition: cg_types.hpp:135
sycl::_V1::detail::ArgDesc::MIndex
int MIndex
Definition: cg_types.hpp:48