Skip to content

Commit 4009536

Browse files
aelovikov-intelaelovikov-intel
authored
[SYCL] Deprecate overloads accepting both sycl::kernel and lambda (#18044)
None of the E2E tests fail in "preview" mode so the functionality hasn't even been tested for awhile. Separate the code into a separate legacy helper to avoid touching in future refactorings for the up-to-date code.
1 parent 37fd1cc commit 4009536

File tree

1 file changed

+76
-40
lines changed

1 file changed

+76
-40
lines changed

‎sycl/include/sycl/handler.hpp

Lines changed: 76 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -1285,7 +1285,7 @@ class __SYCL_EXPORT handler {
12851285
using TransformedArgType = sycl::nd_item<Dims>;
12861286

12871287
wrap_kernel<WrapAs::parallel_for, KernelName, TransformedArgType, Dims>(
1288-
KernelFunc, nullptr /*Kernel*/, Props, ExecutionRange);
1288+
KernelFunc, Props, ExecutionRange);
12891289
}
12901290

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

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

13911390
#ifdef SYCL_LANGUAGE_VERSION
@@ -1601,30 +1600,65 @@ class __SYCL_EXPORT handler {
16011600
WrapAs WrapAsVal, typename KernelName, typename ElementType = void,
16021601
int Dims = 1, bool SetNumWorkGroups = false,
16031602
typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
1604-
typename KernelType, typename MaybeKernelTy, typename... RangeParams>
1605-
void wrap_kernel(const KernelType &KernelFunc, MaybeKernelTy &&MaybeKernel,
1606-
const PropertiesT &Props,
1603+
typename KernelType, typename... RangeParams>
1604+
void wrap_kernel(const KernelType &KernelFunc, const PropertiesT &Props,
16071605
[[maybe_unused]] RangeParams &&...params) {
16081606
// TODO: Properties may change the kernel function, so in order to avoid
16091607
// conflicts they should be included in the name.
16101608
using NameT =
16111609
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
16121610
(void)Props;
1613-
(void)MaybeKernel;
1614-
static_assert(std::is_same_v<MaybeKernelTy, kernel> ||
1615-
std::is_same_v<MaybeKernelTy, std::nullptr_t>);
16161611
KernelWrapper<WrapAsVal, NameT, KernelType, ElementType, PropertiesT>::wrap(
16171612
this, KernelFunc);
16181613
#ifndef __SYCL_DEVICE_ONLY__
16191614
if constexpr (WrapAsVal == WrapAs::single_task) {
16201615
throwOnKernelParameterMisuse<KernelName, KernelType>();
16211616
}
16221617
throwIfActionIsCreated();
1623-
if constexpr (std::is_same_v<MaybeKernelTy, kernel>) {
1624-
// Ignore any set kernel bundles and use the one associated with the
1625-
// kernel.
1626-
setHandlerKernelBundle(MaybeKernel);
1618+
verifyUsedKernelBundleInternal(
1619+
detail::string_view{detail::getKernelName<NameT>()});
1620+
setType(detail::CGType::Kernel);
1621+
1622+
detail::checkValueRange<Dims>(params...);
1623+
if constexpr (SetNumWorkGroups) {
1624+
setNDRangeDescriptor(std::move(params)...,
1625+
/*SetNumWorkGroups=*/true);
1626+
} else {
1627+
setNDRangeDescriptor(std::move(params)...);
16271628
}
1629+
1630+
StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
1631+
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
1632+
#endif
1633+
}
1634+
1635+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
1636+
// Implementation for something that had to be removed long ago but now stuck
1637+
// until next major release...
1638+
template <
1639+
WrapAs WrapAsVal, typename KernelName, typename ElementType = void,
1640+
int Dims = 1, bool SetNumWorkGroups = false,
1641+
typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
1642+
typename KernelType, typename... RangeParams>
1643+
void wrap_kernel_legacy(const KernelType &KernelFunc, kernel &Kernel,
1644+
const PropertiesT &Props,
1645+
[[maybe_unused]] RangeParams &&...params) {
1646+
// TODO: Properties may change the kernel function, so in order to avoid
1647+
// conflicts they should be included in the name.
1648+
using NameT =
1649+
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
1650+
(void)Props;
1651+
(void)Kernel;
1652+
KernelWrapper<WrapAsVal, NameT, KernelType, ElementType, PropertiesT>::wrap(
1653+
this, KernelFunc);
1654+
#ifndef __SYCL_DEVICE_ONLY__
1655+
if constexpr (WrapAsVal == WrapAs::single_task) {
1656+
throwOnKernelParameterMisuse<KernelName, KernelType>();
1657+
}
1658+
throwIfActionIsCreated();
1659+
// Ignore any set kernel bundles and use the one associated with the
1660+
// kernel.
1661+
setHandlerKernelBundle(Kernel);
16281662
verifyUsedKernelBundleInternal(
16291663
detail::string_view{detail::getKernelName<NameT>()});
16301664
setType(detail::CGType::Kernel);
@@ -1637,21 +1671,17 @@ class __SYCL_EXPORT handler {
16371671
setNDRangeDescriptor(std::move(params)...);
16381672
}
16391673

1640-
if constexpr (std::is_same_v<MaybeKernelTy, std::nullptr_t>) {
1641-
StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
1674+
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1675+
if (!lambdaAndKernelHaveEqualName<NameT>()) {
1676+
extractArgsAndReqs();
1677+
MKernelName = getKernelName();
16421678
} else {
1643-
MKernel = detail::getSyclObjImpl(std::move(MaybeKernel));
1644-
if (!lambdaAndKernelHaveEqualName<NameT>()) {
1645-
extractArgsAndReqs();
1646-
MKernelName = getKernelName();
1647-
} else {
1648-
StoreLambda<NameT, KernelType, Dims, ElementType>(
1649-
std::move(KernelFunc));
1650-
}
1679+
StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
16511680
}
16521681
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
16531682
#endif
16541683
}
1684+
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
16551685

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

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

19481977
/// Hierarchical kernel invocation method of a kernel defined as a lambda
@@ -2090,6 +2119,7 @@ class __SYCL_EXPORT handler {
20902119
#endif
20912120
}
20922121

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

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

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

@@ -2161,14 +2193,15 @@ class __SYCL_EXPORT handler {
21612193
/// \param KernelFunc is a lambda representing kernel.
21622194
template <typename KernelName = detail::auto_name, typename KernelType,
21632195
int Dims>
2196+
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
21642197
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
21652198
const KernelType &KernelFunc) {
21662199
using LambdaArgType =
21672200
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
2168-
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
2169-
Dims,
2170-
/*SetNumWorkGroups*/ true>(KernelFunc, Kernel, {} /*Props*/,
2171-
NumWorkGroups);
2201+
wrap_kernel_legacy<WrapAs::parallel_for_work_group, KernelName,
2202+
LambdaArgType, Dims,
2203+
/*SetNumWorkGroups*/ true>(KernelFunc, Kernel,
2204+
{} /*Props*/, NumWorkGroups);
21722205
}
21732206

21742207
/// Hierarchical kernel invocation method of a kernel.
@@ -2188,16 +2221,19 @@ class __SYCL_EXPORT handler {
21882221
/// \param KernelFunc is a lambda representing kernel.
21892222
template <typename KernelName = detail::auto_name, typename KernelType,
21902223
int Dims>
2224+
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
21912225
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
21922226
range<Dims> WorkGroupSize,
21932227
const KernelType &KernelFunc) {
21942228
using LambdaArgType =
21952229
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
21962230
nd_range<Dims> ExecRange =
21972231
nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
2198-
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
2199-
Dims>(KernelFunc, Kernel, {} /*Props*/, ExecRange);
2232+
wrap_kernel_legacy<WrapAs::parallel_for_work_group, KernelName,
2233+
LambdaArgType, Dims>(KernelFunc, Kernel, {} /*Props*/,
2234+
ExecRange);
22002235
}
2236+
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
22012237

22022238
template <typename KernelName = detail::auto_name, typename KernelType,
22032239
typename PropertiesT>
@@ -2208,8 +2244,8 @@ class __SYCL_EXPORT handler {
22082244
std::enable_if_t<ext::oneapi::experimental::is_property_list<
22092245
PropertiesT>::value> single_task(PropertiesT Props,
22102246
const KernelType &KernelFunc) {
2211-
wrap_kernel<WrapAs::single_task, KernelName>(KernelFunc, nullptr /*Kernel*/,
2212-
Props, range<1>{1});
2247+
wrap_kernel<WrapAs::single_task, KernelName>(KernelFunc, Props,
2248+
range<1>{1});
22132249
}
22142250

22152251
template <typename KernelName = detail::auto_name, typename KernelType,

0 commit comments

Comments
 (0)