Skip to content

[SYCL] Deprecate overloads accepting both sycl::kernel and lambda #18044

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

Merged
merged 1 commit into from
Apr 16, 2025
Merged

[SYCL] Deprecate overloads accepting both sycl::kernel and lambda #18044

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
116 changes: 76 additions & 40 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1285,7 +1285,7 @@ class __SYCL_EXPORT handler {
using TransformedArgType = sycl::nd_item<Dims>;

wrap_kernel<WrapAs::parallel_for, KernelName, TransformedArgType, Dims>(
KernelFunc, nullptr /*Kernel*/, Props, ExecutionRange);
KernelFunc, Props, ExecutionRange);
}

/// Defines and invokes a SYCL kernel function for the specified range.
Expand Down Expand Up @@ -1357,8 +1357,7 @@ class __SYCL_EXPORT handler {
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims,
/*SetNumWorkGroups=*/true>(KernelFunc, nullptr /*Kernel*/,
Props, NumWorkGroups);
/*SetNumWorkGroups=*/true>(KernelFunc, Props, NumWorkGroups);
}

/// Hierarchical kernel invocation method of a kernel defined as a lambda
Expand All @@ -1385,7 +1384,7 @@ class __SYCL_EXPORT handler {
nd_range<Dims> ExecRange =
nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims>(KernelFunc, nullptr /*Kernel*/, Props, ExecRange);
Dims>(KernelFunc, Props, ExecRange);
}

#ifdef SYCL_LANGUAGE_VERSION
Expand Down Expand Up @@ -1601,30 +1600,65 @@ class __SYCL_EXPORT handler {
WrapAs WrapAsVal, typename KernelName, typename ElementType = void,
int Dims = 1, bool SetNumWorkGroups = false,
typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
typename KernelType, typename MaybeKernelTy, typename... RangeParams>
void wrap_kernel(const KernelType &KernelFunc, MaybeKernelTy &&MaybeKernel,
const PropertiesT &Props,
typename KernelType, typename... RangeParams>
void wrap_kernel(const KernelType &KernelFunc, const PropertiesT &Props,
[[maybe_unused]] RangeParams &&...params) {
// TODO: Properties may change the kernel function, so in order to avoid
// conflicts they should be included in the name.
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
(void)Props;
(void)MaybeKernel;
static_assert(std::is_same_v<MaybeKernelTy, kernel> ||
std::is_same_v<MaybeKernelTy, std::nullptr_t>);
KernelWrapper<WrapAsVal, NameT, KernelType, ElementType, PropertiesT>::wrap(
this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
if constexpr (WrapAsVal == WrapAs::single_task) {
throwOnKernelParameterMisuse<KernelName, KernelType>();
}
throwIfActionIsCreated();
if constexpr (std::is_same_v<MaybeKernelTy, kernel>) {
// Ignore any set kernel bundles and use the one associated with the
// kernel.
setHandlerKernelBundle(MaybeKernel);
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
setType(detail::CGType::Kernel);

detail::checkValueRange<Dims>(params...);
if constexpr (SetNumWorkGroups) {
setNDRangeDescriptor(std::move(params)...,
/*SetNumWorkGroups=*/true);
} else {
setNDRangeDescriptor(std::move(params)...);
}

StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
#endif
}

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// Implementation for something that had to be removed long ago but now stuck
// until next major release...
template <
WrapAs WrapAsVal, typename KernelName, typename ElementType = void,
int Dims = 1, bool SetNumWorkGroups = false,
typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
typename KernelType, typename... RangeParams>
void wrap_kernel_legacy(const KernelType &KernelFunc, kernel &Kernel,
const PropertiesT &Props,
[[maybe_unused]] RangeParams &&...params) {
// TODO: Properties may change the kernel function, so in order to avoid
// conflicts they should be included in the name.
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
(void)Props;
(void)Kernel;
KernelWrapper<WrapAsVal, NameT, KernelType, ElementType, PropertiesT>::wrap(
this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
if constexpr (WrapAsVal == WrapAs::single_task) {
throwOnKernelParameterMisuse<KernelName, KernelType>();
}
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the
// kernel.
setHandlerKernelBundle(Kernel);
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
setType(detail::CGType::Kernel);
Expand All @@ -1637,21 +1671,17 @@ class __SYCL_EXPORT handler {
setNDRangeDescriptor(std::move(params)...);
}

if constexpr (std::is_same_v<MaybeKernelTy, std::nullptr_t>) {
StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
if (!lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
MKernelName = getKernelName();
} else {
MKernel = detail::getSyclObjImpl(std::move(MaybeKernel));
if (!lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
MKernelName = getKernelName();
} else {
StoreLambda<NameT, KernelType, Dims, ElementType>(
std::move(KernelFunc));
}
StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
}
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
#endif
}
#endif // __INTEL_PREVIEW_BREAKING_CHANGES

// NOTE: to support kernel_handler argument in kernel lambdas, only
// KernelWrapper<...>::wrap() must be called in this code.
Expand Down Expand Up @@ -1873,8 +1903,8 @@ class __SYCL_EXPORT handler {
/// \param KernelFunc is a SYCL kernel function.
template <typename KernelName = detail::auto_name, typename KernelType>
void single_task(const KernelType &KernelFunc) {
wrap_kernel<WrapAs::single_task, KernelName>(KernelFunc, nullptr /*Kernel*/,
{} /*Props*/, range<1>{1});
wrap_kernel<WrapAs::single_task, KernelName>(KernelFunc, {} /*Props*/,
range<1>{1});
}

template <typename KernelName = detail::auto_name, typename KernelType>
Expand Down Expand Up @@ -1941,8 +1971,7 @@ class __SYCL_EXPORT handler {
std::is_integral<LambdaArgType>::value && Dims == 1, item<Dims>,
typename TransformUserItemType<Dims, LambdaArgType>::type>;
wrap_kernel<WrapAs::parallel_for, KernelName, TransformedArgType, Dims>(
KernelFunc, nullptr /*Kernel*/, {} /*Props*/, NumWorkItems,
WorkItemOffset);
KernelFunc, {} /*Props*/, NumWorkItems, WorkItemOffset);
}

/// Hierarchical kernel invocation method of a kernel defined as a lambda
Expand Down Expand Up @@ -2090,6 +2119,7 @@ class __SYCL_EXPORT handler {
#endif
}

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
/// Defines and invokes a SYCL kernel function for the specified range.
///
/// \param Kernel is a SYCL kernel that is executed on a SYCL device
Expand All @@ -2099,12 +2129,13 @@ class __SYCL_EXPORT handler {
/// is a host device.
template <typename KernelName = detail::auto_name, typename KernelType,
int Dims>
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
const KernelType &KernelFunc) {
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(Kernel);
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
wrap_kernel<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
wrap_kernel_legacy<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
KernelFunc, Kernel, {} /*Props*/, NumWorkItems);
}

Expand All @@ -2119,11 +2150,11 @@ class __SYCL_EXPORT handler {
/// is a host device.
template <typename KernelName = detail::auto_name, typename KernelType,
int Dims>
__SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
id<Dims> WorkItemOffset, const KernelType &KernelFunc) {
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
wrap_kernel<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
wrap_kernel_legacy<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
KernelFunc, Kernel, {} /*Props*/, NumWorkItems, WorkItemOffset);
}

Expand All @@ -2138,11 +2169,12 @@ class __SYCL_EXPORT handler {
/// is a host device.
template <typename KernelName = detail::auto_name, typename KernelType,
int Dims>
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
void parallel_for(kernel Kernel, nd_range<Dims> NDRange,
const KernelType &KernelFunc) {
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
wrap_kernel<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
wrap_kernel_legacy<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
KernelFunc, Kernel, {} /*Props*/, NDRange);
}

Expand All @@ -2161,14 +2193,15 @@ class __SYCL_EXPORT handler {
/// \param KernelFunc is a lambda representing kernel.
template <typename KernelName = detail::auto_name, typename KernelType,
int Dims>
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
const KernelType &KernelFunc) {
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims,
/*SetNumWorkGroups*/ true>(KernelFunc, Kernel, {} /*Props*/,
NumWorkGroups);
wrap_kernel_legacy<WrapAs::parallel_for_work_group, KernelName,
LambdaArgType, Dims,
/*SetNumWorkGroups*/ true>(KernelFunc, Kernel,
{} /*Props*/, NumWorkGroups);
}

/// Hierarchical kernel invocation method of a kernel.
Expand All @@ -2188,16 +2221,19 @@ class __SYCL_EXPORT handler {
/// \param KernelFunc is a lambda representing kernel.
template <typename KernelName = detail::auto_name, typename KernelType,
int Dims>
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
range<Dims> WorkGroupSize,
const KernelType &KernelFunc) {
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
nd_range<Dims> ExecRange =
nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims>(KernelFunc, Kernel, {} /*Props*/, ExecRange);
wrap_kernel_legacy<WrapAs::parallel_for_work_group, KernelName,
LambdaArgType, Dims>(KernelFunc, Kernel, {} /*Props*/,
ExecRange);
}
#endif // __INTEL_PREVIEW_BREAKING_CHANGES

template <typename KernelName = detail::auto_name, typename KernelType,
typename PropertiesT>
Expand All @@ -2208,8 +2244,8 @@ class __SYCL_EXPORT handler {
std::enable_if_t<ext::oneapi::experimental::is_property_list<
PropertiesT>::value> single_task(PropertiesT Props,
const KernelType &KernelFunc) {
wrap_kernel<WrapAs::single_task, KernelName>(KernelFunc, nullptr /*Kernel*/,
Props, range<1>{1});
wrap_kernel<WrapAs::single_task, KernelName>(KernelFunc, Props,
range<1>{1});
}

template <typename KernelName = detail::auto_name, typename KernelType,
Expand Down