31 #include <type_traits>
35 inline namespace _V1 {
56 void setNDRangeLeftover(
int Dims_) {
57 for (
int I = Dims_; I < 3; ++I) {
70 for (
int I = 0; I < Dims_; ++I) {
76 setNDRangeLeftover(Dims_);
84 for (
int I = 0; I < Dims_; ++I) {
90 setNDRangeLeftover(Dims_);
95 for (
int I = 0; I < Dims_; ++I) {
101 setNDRangeLeftover(Dims_);
106 for (
int I = 0; I < Dims_; ++I) {
112 setNDRangeLeftover(Dims_);
117 for (
int I = 0; I < Dims_; ++I) {
124 setNDRangeLeftover(Dims_);
139 static_assert(std::integral_constant<T, false>::value,
140 "Second template parameter is required to be of function type");
143 template <
typename F,
typename RetT,
typename... Args>
146 template <
typename T>
147 static constexpr
auto check(T *) ->
typename std::is_same<
148 decltype(std::declval<T>().
operator()(std::declval<Args>()...)),
151 template <
typename>
static constexpr std::false_type check(...);
153 using type = decltype(check<F>(0));
156 static constexpr
bool value = type::value;
159 template <
typename F,
typename... Args>
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>();
177 typename KernelType,
typename LambdaArgType,
178 typename std::enable_if_t<!std::is_same_v<LambdaArgType, void>> * =
nullptr>
185 template <
typename KernelType,
typename LambdaArgType =
void>
188 isKernelLambdaCallableWithKernelHandlerImpl<KernelType, LambdaArgType>();
193 template <
typename KernelType>
194 typename std::enable_if_t<KernelLambdaHasKernelHandlerArgT<KernelType>::value>
200 template <
typename KernelType>
201 typename std::enable_if_t<!KernelLambdaHasKernelHandlerArgT<KernelType>::value>
206 template <
typename ArgType,
typename KernelType>
207 typename std::enable_if_t<
214 template <
typename ArgType,
typename KernelType>
215 typename std::enable_if_t<
228 virtual char *
getPtr() = 0;
233 std::function<void()> MHostTask;
238 HostTask(std::function<
void()> &&Func) : MHostTask(Func) {}
248 template <
class KernelType,
class KernelArgType,
int Dims>
250 using IDBuilder = sycl::detail::Builder;
270 AdjustedRange.
set(NDRDesc.
Dims,
276 for (
size_t I = 0; I < AdjustedRange.
Dims; ++I)
285 char *
getPtr()
override {
return reinterpret_cast<char *
>(&MKernel); }
287 template <
class ArgT = KernelArgType>
288 typename std::enable_if_t<std::is_same_v<ArgT, void>>
293 template <
class ArgT = KernelArgType>
294 typename std::enable_if_t<std::is_same_v<ArgT, sycl::id<Dims>>>
302 for (
int I = 0; I < Dims; ++I) {
305 UpperBound[I] = Range[I] + Offset[I];
309 Offset, Stride, UpperBound,
312 IDBuilder::createItem<Dims, true>(Range, ID, Offset);
314 runKernelWithArg<const sycl::id<Dims> &>(MKernel, ID);
318 template <
class ArgT = KernelArgType>
319 typename std::enable_if_t<std::is_same_v<ArgT,
item<Dims,
false>>>
323 for (
int I = 0; I < Dims; ++I)
327 sycl::item<Dims, false> Item =
328 IDBuilder::createItem<Dims, false>(Range, ID);
329 sycl::item<Dims, true> ItemWithOffset = Item;
331 runKernelWithArg<sycl::item<Dims, false>>(MKernel, Item);
335 template <
class ArgT = KernelArgType>
336 typename std::enable_if_t<std::is_same_v<ArgT,
item<Dims,
true>>>
344 for (
int I = 0; I < Dims; ++I) {
347 UpperBound[I] = Range[I] + Offset[I];
351 Offset, Stride, UpperBound,
354 IDBuilder::createItem<Dims, true>(Range, ID, Offset);
360 template <
class ArgT = KernelArgType>
361 typename std::enable_if_t<std::is_same_v<ArgT, nd_item<Dims>>>
364 for (
int I = 0; I < Dims; ++I) {
367 throw sycl::nd_range_error(
"Invalid local size for global size",
368 PI_ERROR_INVALID_WORK_GROUP_SIZE);
376 for (
int I = 0; I < Dims; ++I) {
384 GlobalSize, LocalSize, GroupSize, GroupID);
387 id<Dims> GlobalID = GroupID * LocalSize + LocalID + GlobalOffset;
389 IDBuilder::createItem<Dims, true>(GlobalSize, GlobalID,
392 IDBuilder::createItem<Dims, false>(LocalSize, LocalID);
394 IDBuilder::createNDItem<Dims>(GlobalItem, LocalItem, Group);
396 runKernelWithArg<const sycl::nd_item<Dims>>(MKernel, NDItem);
401 template <
typename ArgT = KernelArgType>
402 std::enable_if_t<std::is_same_v<ArgT, sycl::group<Dims>>>
406 for (
int I = 0; I < Dims; ++I) {
409 throw sycl::nd_range_error(
"Invalid local size for global size",
410 PI_ERROR_INVALID_WORK_GROUP_SIZE);
417 for (
int I = 0; I < Dims; ++I) {
423 IDBuilder::createGroup<Dims>(GlobalSize, LocalSize, NGroups, GroupID);
424 runKernelWithArg<sycl::group<Dims>>(MKernel, Group);