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