Skip to content

[SYCL][NFCI] More refactoring around "kernel wrapping" #18015

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 14, 2025
Merged

[SYCL][NFCI] More refactoring around "kernel wrapping" #18015

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
272 changes: 78 additions & 194 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1267,10 +1267,6 @@ class __SYCL_EXPORT handler {
typename PropertiesT>
void parallel_for_impl(nd_range<Dims> ExecutionRange, PropertiesT Props,
const KernelType &KernelFunc) {
// 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;
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
static_assert(
Expand All @@ -1279,21 +1275,8 @@ class __SYCL_EXPORT handler {
"must be either sycl::nd_item or be convertible from sycl::nd_item");
using TransformedArgType = sycl::nd_item<Dims>;

(void)ExecutionRange;
(void)Props;
KernelWrapper<WrapAs::parallel_for, NameT, KernelType, TransformedArgType,
PropertiesT>::wrap(this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
detail::checkValueRange<Dims>(ExecutionRange);
setNDRangeDescriptor(std::move(ExecutionRange));
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
std::move(KernelFunc));
setType(detail::CGType::Kernel);
#endif
wrap_kernel<WrapAs::parallel_for, KernelName, TransformedArgType, Dims>(
KernelFunc, nullptr /*Kernel*/, Props, ExecutionRange);
}

/// Defines and invokes a SYCL kernel function for the specified range.
Expand Down Expand Up @@ -1361,26 +1344,12 @@ class __SYCL_EXPORT handler {
void parallel_for_work_group_lambda_impl(range<Dims> NumWorkGroups,
PropertiesT Props,
const KernelType &KernelFunc) {
// 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;
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
(void)NumWorkGroups;
(void)Props;
KernelWrapper<WrapAs::parallel_for_work_group, NameT, KernelType,
LambdaArgType, PropertiesT>::wrap(this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
detail::checkValueRange<Dims>(NumWorkGroups);
setNDRangeDescriptor(NumWorkGroups, /*SetNumWorkGroups=*/true);
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
setType(detail::CGType::Kernel);
#endif // __SYCL_DEVICE_ONLY__
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims,
/*SetNumWorkGroups=*/true>(KernelFunc, nullptr /*Kernel*/,
Props, NumWorkGroups);
}

/// Hierarchical kernel invocation method of a kernel defined as a lambda
Expand All @@ -1402,29 +1371,12 @@ class __SYCL_EXPORT handler {
range<Dims> WorkGroupSize,
PropertiesT Props,
const KernelType &KernelFunc) {
// 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;
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
(void)NumWorkGroups;
(void)WorkGroupSize;
(void)Props;
KernelWrapper<WrapAs::parallel_for_work_group, NameT, KernelType,
LambdaArgType, PropertiesT>::wrap(this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
nd_range<Dims> ExecRange =
nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
detail::checkValueRange<Dims>(ExecRange);
setNDRangeDescriptor(std::move(ExecRange));
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
setType(detail::CGType::Kernel);
#endif // __SYCL_DEVICE_ONLY__
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims>(KernelFunc, nullptr /*Kernel*/, Props, ExecRange);
}

#ifdef SYCL_LANGUAGE_VERSION
Expand Down Expand Up @@ -1636,6 +1588,59 @@ class __SYCL_EXPORT handler {
}
};

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 MaybeKernelTy, typename... RangeParams>
void wrap_kernel(const KernelType &KernelFunc, MaybeKernelTy &&MaybeKernel,
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__
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)...);
}

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

// NOTE: to support kernel_handler argument in kernel lambdas, only
// KernelWrapper<...>::wrap() must be called in this code.

Expand All @@ -1651,25 +1656,10 @@ class __SYCL_EXPORT handler {
typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
void single_task_lambda_impl(PropertiesT Props,
const KernelType &KernelFunc) {
(void)Props;
// 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;

KernelWrapper<WrapAs::single_task, NameT, KernelType, void,
PropertiesT>::wrap(this, KernelFunc);
wrap_kernel<WrapAs::single_task, KernelName>(KernelFunc, nullptr /*Kernel*/,
Props, range<1>{1});
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
throwOnKernelParameterMisuse<KernelName, KernelType>();
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
// No need to check if range is out of INT_MAX limits as it's compile-time
// known constant.
setNDRangeDescriptor(range<1>{1});
processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
StoreLambda<NameT, KernelType, /*Dims*/ 1, void>(KernelFunc);
setType(detail::CGType::Kernel);
#endif
}

Expand Down Expand Up @@ -1953,26 +1943,13 @@ class __SYCL_EXPORT handler {
__SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020")
void parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
const KernelType &KernelFunc) {
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
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>;
(void)NumWorkItems;
(void)WorkItemOffset;
KernelWrapper<WrapAs::parallel_for, NameT, KernelType,
TransformedArgType>::wrap(this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
setNDRangeDescriptor(std::move(NumWorkItems), std::move(WorkItemOffset));
StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
std::move(KernelFunc));
setType(detail::CGType::Kernel);
#endif
wrap_kernel<WrapAs::parallel_for, KernelName, TransformedArgType, Dims>(
KernelFunc, nullptr /*Kernel*/, {} /*Props*/, NumWorkItems,
WorkItemOffset);
}

/// Hierarchical kernel invocation method of a kernel defined as a lambda
Expand Down Expand Up @@ -2133,28 +2110,9 @@ class __SYCL_EXPORT handler {
const KernelType &KernelFunc) {
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(Kernel);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
(void)Kernel;
(void)NumWorkItems;
KernelWrapper<WrapAs::parallel_for, NameT, KernelType, LambdaArgType>::wrap(
this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
detail::checkValueRange<Dims>(NumWorkItems);
setNDRangeDescriptor(std::move(NumWorkItems));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
setType(detail::CGType::Kernel);
if (!lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
MKernelName = getKernelName();
} else
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
std::move(KernelFunc));
#endif
wrap_kernel<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
KernelFunc, Kernel, {} /*Props*/, NumWorkItems);
}

/// Defines and invokes a SYCL kernel function for the specified range and
Expand All @@ -2171,31 +2129,9 @@ class __SYCL_EXPORT handler {
__SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
id<Dims> WorkItemOffset, const KernelType &KernelFunc) {
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
(void)Kernel;
(void)NumWorkItems;
(void)WorkItemOffset;
KernelWrapper<WrapAs::parallel_for, NameT, KernelType, LambdaArgType>::wrap(
this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(Kernel);
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
setNDRangeDescriptor(std::move(NumWorkItems), std::move(WorkItemOffset));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
setType(detail::CGType::Kernel);
if (!lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
MKernelName = getKernelName();
} else
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
std::move(KernelFunc));
#endif
wrap_kernel<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
KernelFunc, Kernel, {} /*Props*/, NumWorkItems, WorkItemOffset);
}

/// Defines and invokes a SYCL kernel function for the specified range and
Expand All @@ -2211,31 +2147,10 @@ class __SYCL_EXPORT handler {
int Dims>
void parallel_for(kernel Kernel, nd_range<Dims> NDRange,
const KernelType &KernelFunc) {
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
(void)Kernel;
(void)NDRange;
KernelWrapper<WrapAs::parallel_for, NameT, KernelType, LambdaArgType>::wrap(
this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(Kernel);
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
detail::checkValueRange<Dims>(NDRange);
setNDRangeDescriptor(std::move(NDRange));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
setType(detail::CGType::Kernel);
if (!lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
MKernelName = getKernelName();
} else
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
std::move(KernelFunc));
#endif
wrap_kernel<WrapAs::parallel_for, KernelName, LambdaArgType, Dims>(
KernelFunc, Kernel, {} /*Props*/, NDRange);
}

/// Hierarchical kernel invocation method of a kernel.
Expand All @@ -2255,26 +2170,12 @@ class __SYCL_EXPORT handler {
int Dims>
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
const KernelType &KernelFunc) {
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
(void)Kernel;
(void)NumWorkGroups;
KernelWrapper<WrapAs::parallel_for_work_group, NameT, KernelType,
LambdaArgType>::wrap(this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(Kernel);
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
detail::checkValueRange<Dims>(NumWorkGroups);
setNDRangeDescriptor(NumWorkGroups, /*SetNumWorkGroups=*/true);
MKernel = detail::getSyclObjImpl(std::move(Kernel));
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
setType(detail::CGType::Kernel);
#endif // __SYCL_DEVICE_ONLY__
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims,
/*SetNumWorkGroups*/ true>(KernelFunc, Kernel, {} /*Props*/,
NumWorkGroups);
}

/// Hierarchical kernel invocation method of a kernel.
Expand All @@ -2297,29 +2198,12 @@ class __SYCL_EXPORT handler {
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
range<Dims> WorkGroupSize,
const KernelType &KernelFunc) {
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
using LambdaArgType =
sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
(void)Kernel;
(void)NumWorkGroups;
(void)WorkGroupSize;
KernelWrapper<WrapAs::parallel_for_work_group, NameT, KernelType,
LambdaArgType>::wrap(this, KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(Kernel);
verifyUsedKernelBundleInternal(
detail::string_view{detail::getKernelName<NameT>()});
nd_range<Dims> ExecRange =
nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
detail::checkValueRange<Dims>(ExecRange);
setNDRangeDescriptor(std::move(ExecRange));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
setType(detail::CGType::Kernel);
#endif // __SYCL_DEVICE_ONLY__
wrap_kernel<WrapAs::parallel_for_work_group, KernelName, LambdaArgType,
Dims>(KernelFunc, Kernel, {} /*Props*/, ExecRange);
}

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