Skip to content

[SYCL] Handler-less kernel submit path (range based) #20912

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
slawekptak wants to merge 4 commits into
base: sycl
Choose a base branch
from
+446 −174
Open

[SYCL] Handler-less kernel submit path (range based) #20912

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b3e6ff9
Handler-less kernel submit - range-based
slawekptak Dec 15, 2025
eb80bf1
Update include path
slawekptak Dec 16, 2025
be11e1e
Use std::decay_t for KernelType
slawekptak Dec 16, 2025
85bd4bd
Merge branch 'sycl' into no_handler_range_v2
slawekptak Dec 17, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions sycl/include/sycl/detail/nd_range_view.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,10 @@ class nd_range_view {
nd_range_view &operator=(const nd_range_view &Desc) = default;
nd_range_view &operator=(nd_range_view &&Desc) = default;

template <int Dims_>
Copy link
Contributor

@cperkinsintel cperkinsintel Dec 19, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

shouldn't these be templated with <size_t Dims_> instead of int ? Then we won't need the cast.

Copy link
Contributor

@uditagarwal97 uditagarwal97 Dec 19, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

+1. And also nd_range_view implies view of sycl::nd_range, not sycl::range. I think we should have an anologous range_view as well.

nd_range_view(sycl::range<Dims_> &N)
: MGlobalSize(&(N[0])), MDims(size_t(Dims_)) {}

template <int Dims_>
nd_range_view(sycl::nd_range<Dims_> &ExecutionRange)
: MGlobalSize(&(ExecutionRange.globalSize[0])),
Expand Down
53 changes: 45 additions & 8 deletions sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -192,10 +192,26 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
typename KernelType, typename... ReductionsT>
void parallel_for(queue Q, range<Dimensions> Range, const KernelType &KernelObj,
ReductionsT &&...Reductions) {
submit(std::move(Q), [&](handler &CGH) {
parallel_for<KernelName>(CGH, Range, KernelObj,
std::forward<ReductionsT>(Reductions)...);
});
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, item<Dimensions>>;
using TransformedArgType = std::conditional_t<
std::is_integral<LambdaArgType>::value && Dimensions == 1,
item<Dimensions>,
typename detail::TransformUserItemType<Dimensions, LambdaArgType>::type>;

// TODO The handler-less path does not support reductions, and
// kernel functions with the kernel_handler type argument yet.
if constexpr (sizeof...(ReductionsT) == 0 &&
!(detail::KernelLambdaHasKernelHandlerArgT<
KernelType, TransformedArgType>::value)) {
detail::submit_kernel_direct_parallel_for<KernelName>(std::move(Q), Range,
KernelObj);
} else {
submit(std::move(Q), [&](handler &CGH) {
parallel_for<KernelName>(CGH, Range, KernelObj,
std::forward<ReductionsT>(Reductions)...);
});
}
}

template <typename KernelName = sycl::detail::auto_name, int Dimensions,
Expand All @@ -215,10 +231,31 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
typename Properties, typename KernelType, typename... ReductionsT>
void parallel_for(queue Q, launch_config<range<Dimensions>, Properties> Config,
const KernelType &KernelObj, ReductionsT &&...Reductions) {
submit(std::move(Q), [&](handler &CGH) {
parallel_for<KernelName>(CGH, Config, KernelObj,
std::forward<ReductionsT>(Reductions)...);
});
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, item<Dimensions>>;
using TransformedArgType = std::conditional_t<
std::is_integral<LambdaArgType>::value && Dimensions == 1,
item<Dimensions>,
typename detail::TransformUserItemType<Dimensions, LambdaArgType>::type>;

// TODO The handler-less path does not support reductions, and
// kernel functions with the kernel_handler type argument yet.
if constexpr (sizeof...(ReductionsT) == 0 &&
!(detail::KernelLambdaHasKernelHandlerArgT<
KernelType, TransformedArgType>::value)) {
ext::oneapi::experimental::detail::LaunchConfigAccess<range<Dimensions>,
Properties>
LaunchConfigAccess(Config);

detail::submit_kernel_direct_parallel_for<KernelName>(
std::move(Q), LaunchConfigAccess.getRange(), KernelObj, {},
LaunchConfigAccess.getProperties());
} else {
submit(std::move(Q), [&](handler &CGH) {
parallel_for<KernelName>(CGH, Config, KernelObj,
std::forward<ReductionsT>(Reductions)...);
});
}
}

template <int Dimensions, typename... ArgsT>
Expand Down
18 changes: 9 additions & 9 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -279,6 +279,13 @@ template <int Dims> bool range_size_fits_in_size_t(const range<Dims> &r) {
return true;
}

template <int Dims, typename LambdaArgType> struct TransformUserItemType {
using type = std::conditional_t<
std::is_convertible_v<nd_item<Dims>, LambdaArgType>, nd_item<Dims>,
std::conditional_t<std::is_convertible_v<item<Dims>, LambdaArgType>,
item<Dims>, LambdaArgType>>;
};

} // namespace detail

/// Command group handler class.
Expand Down Expand Up @@ -778,13 +785,6 @@ class __SYCL_EXPORT handler {

device get_device() const;

template <int Dims, typename LambdaArgType> struct TransformUserItemType {
using type = std::conditional_t<
std::is_convertible_v<nd_item<Dims>, LambdaArgType>, nd_item<Dims>,
std::conditional_t<std::is_convertible_v<item<Dims>, LambdaArgType>,
item<Dims>, LambdaArgType>>;
};

/// Defines and invokes a SYCL kernel function for the specified range.
///
/// The SYCL kernel function is defined as a lambda function or a named
Expand Down Expand Up @@ -823,7 +823,7 @@ class __SYCL_EXPORT handler {
// sycl::item/sycl::nd_item to transport item information
using TransformedArgType = std::conditional_t<
std::is_integral<LambdaArgType>::value && Dims == 1, item<Dims>,
typename TransformUserItemType<Dims, LambdaArgType>::type>;
typename detail::TransformUserItemType<Dims, LambdaArgType>::type>;

static_assert(!std::is_same_v<TransformedArgType, sycl::nd_item<Dims>>,
"Kernel argument cannot have a sycl::nd_item type in "
Expand Down Expand Up @@ -1349,7 +1349,7 @@ class __SYCL_EXPORT handler {
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
using TransformedArgType = std::conditional_t<
std::is_integral<LambdaArgType>::value && Dims == 1, item<Dims>,
typename TransformUserItemType<Dims, LambdaArgType>::type>;
typename detail::TransformUserItemType<Dims, LambdaArgType>::type>;
wrap_kernel<detail::WrapAs::parallel_for, KernelName, TransformedArgType,
Dims>(KernelFunc, {} /*Props*/, NumWorkItems, WorkItemOffset);
}
Expand Down
40 changes: 37 additions & 3 deletions sycl/include/sycl/khr/free_function_commands.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -54,19 +54,53 @@ template <typename KernelType>
void launch(const queue &q, range<1> r, const KernelType &k,
const sycl::detail::code_location &codeLoc =
sycl::detail::code_location::current()) {
submit(q, [&](handler &h) { launch<KernelType>(h, r, k); }, codeLoc);
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<1>>;
using TransformedArgType = std::conditional_t<
std::is_integral<LambdaArgType>::value, item<1>,
typename detail::TransformUserItemType<1, LambdaArgType>::type>;

// TODO The handler-less path does not support kernel functions with the
// kernel_handler type argument yet.
if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<
KernelType, TransformedArgType>::value)) {
detail::submit_kernel_direct_parallel_for(q, r, k);
} else {
submit(q, [&](handler &h) { launch<KernelType>(h, r, k); }, codeLoc);
}
}
template <typename KernelType>
void launch(const queue &q, range<2> r, const KernelType &k,
const sycl::detail::code_location &codeLoc =
sycl::detail::code_location::current()) {
submit(q, [&](handler &h) { launch<KernelType>(h, r, k); }, codeLoc);
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<2>>;
using TransformedArgType =
typename detail::TransformUserItemType<2, LambdaArgType>::type;

// TODO The handler-less path does not support kernel functions with the
// kernel_handler type argument yet.
if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<
KernelType, TransformedArgType>::value)) {
detail::submit_kernel_direct_parallel_for(q, r, k);
} else {
submit(q, [&](handler &h) { launch<KernelType>(h, r, k); }, codeLoc);
}
}
template <typename KernelType>
void launch(const queue &q, range<3> r, const KernelType &k,
const sycl::detail::code_location &codeLoc =
sycl::detail::code_location::current()) {
submit(q, [&](handler &h) { launch<KernelType>(h, r, k); }, codeLoc);
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<3>>;
using TransformedArgType =
typename detail::TransformUserItemType<3, LambdaArgType>::type;

// TODO The handler-less path does not support kernel functions with the
// kernel_handler type argument yet.
if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<
KernelType, TransformedArgType>::value)) {
detail::submit_kernel_direct_parallel_for(q, r, k);
} else {
submit(q, [&](handler &h) { launch<KernelType>(h, r, k); }, codeLoc);
}
}

template <typename... ArgsT>
Expand Down
Loading