Skip to content

[SYCL][ABI-Break] Drop old enqueue functions from handler #20744

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 3 commits into from
Nov 25, 2025
+0 −246
Merged

[SYCL][ABI-Break] Drop old enqueue functions from handler #20744

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
1078abe
[SYCL][ABI-Break] Drop old enqueue functions from handler
sergey-semenov Nov 24, 2025
69c72e3
Update Windows ABI
sergey-semenov Nov 25, 2025
6266716
Merge branch 'sycl' into handlerkernelnameabi
sergey-semenov Nov 25, 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
235 changes: 0 additions & 235 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -471,27 +471,8 @@ class __SYCL_EXPORT handler {
/// Extracts and prepares kernel arguments set via set_arg(s).
void extractArgsAndReqs();

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
/// \return a string containing name of SYCL kernel.
detail::ABINeutralKernelNameStrT getKernelName();

template <typename LambdaNameT> bool lambdaAndKernelHaveEqualName() {
// TODO It is unclear a kernel and a lambda/functor must to be equal or not
// for parallel_for with sycl::kernel and lambda/functor together
// Now if they are equal we extract arguments from lambda/functor for the
// kernel. Else it is necessary use set_atg(s) for resolve the order and
// values of arguments for the kernel.
assert(MKernel && "MKernel is not initialized");
constexpr std::string_view LambdaName =
detail::CompileTimeKernelInfo<LambdaNameT>.Name;
detail::ABINeutralKernelNameStrT KernelName = getKernelName();
return KernelName == LambdaName;
}
#endif

/// Saves the location of user's code passed in \p CodeLoc for future usage in
/// finalize() method.

void saveCodeLoc(detail::code_location CodeLoc, bool IsTopCodeLoc);
void copyCodeLoc(const handler &other);

Expand Down Expand Up @@ -1248,64 +1229,6 @@ class __SYCL_EXPORT handler {
#endif
}

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// Implementation for something that had to be removed long ago but now stuck
// until next major release...
template <
detail::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;
detail::KernelWrapper<WrapAsVal, NameT, KernelType, ElementType,
PropertiesT>::wrap(KernelFunc);

if constexpr (ext::oneapi::experimental::detail::
HasKernelPropertiesGetMethod<const KernelType &>::value) {
SetKernelLaunchpropertiesIfNotEmpty(detail::extractKernelProperties(
KernelFunc.get(ext::oneapi::experimental::properties_tag{})));
}

#ifndef __SYCL_DEVICE_ONLY__
constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
if constexpr (WrapAsVal == detail::WrapAs::single_task) {
throwOnKernelParameterMisuse(Info);
}
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the
// kernel.
setHandlerKernelBundle(Kernel);
verifyUsedKernelBundleInternal(Info.Name);

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

MKernel = detail::getSyclObjImpl(Kernel);
if (!lambdaAndKernelHaveEqualName<NameT>()) {
throw sycl::exception(
make_error_code(errc::invalid),
"the kernel name must match the name of the lambda");
}
StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
SetKernelLaunchpropertiesIfNotEmpty(
detail::extractKernelProperties<Info.IsESIMD>(Props));
#endif
}
#endif // __INTEL_PREVIEW_BREAKING_CHANGES

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

Expand Down Expand Up @@ -1723,164 +1646,6 @@ class __SYCL_EXPORT handler {
Kernel);
}

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// Implementation for something that had to be removed long ago but now stuck
// until next major release...

/// Defines and invokes a SYCL kernel function.
///
/// \param Kernel is a SYCL kernel that is executed on a SYCL device
/// (except for the host device).
/// \param KernelFunc is a lambda that is used if device, queue is bound to,
/// is a host device.
template <typename KernelName = detail::auto_name, typename KernelType>
__SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.")
void single_task(kernel Kernel, 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;
(void)Kernel;
detail::KernelWrapperHelperFuncs::kernel_single_task<NameT>(KernelFunc);
#ifndef __SYCL_DEVICE_ONLY__
throwIfActionIsCreated();
constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
verifyUsedKernelBundleInternal(Info.Name);
// No need to check if range is out of INT_MAX limits as it's compile-time
// known constant
setNDRangeDescriptor(range<1>{1});

MKernel = detail::getSyclObjImpl(Kernel);
if (!lambdaAndKernelHaveEqualName<NameT>()) {
throw sycl::exception(
make_error_code(errc::invalid),
"the kernel name must match the name of the lambda");
}
StoreLambda<NameT, KernelType, /*Dims*/ 1, void>(std::move(KernelFunc));

#else
detail::CheckDeviceCopyable<KernelType>();
#endif
}
#endif // __INTEL_PREVIEW_BREAKING_CHANGES

#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
/// (except for the host device).
/// \param NumWorkItems is a range defining indexing space.
/// \param KernelFunc is a lambda that is used if device, queue is bound to,
/// 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_legacy<detail::WrapAs::parallel_for, KernelName, LambdaArgType,
Dims>(KernelFunc, Kernel, {} /*Props*/, NumWorkItems);
}

/// Defines and invokes a SYCL kernel function for the specified range and
/// offsets.
///
/// \param Kernel is a SYCL kernel that is executed on a SYCL device
/// (except for the host device).
/// \param NumWorkItems is a range defining indexing space.
/// \param WorkItemOffset is an offset to be applied to each work item index.
/// \param KernelFunc is a lambda that is used if device, queue is bound to,
/// 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,
id<Dims> WorkItemOffset, const KernelType &KernelFunc) {
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
wrap_kernel_legacy<detail::WrapAs::parallel_for, KernelName, LambdaArgType,
Dims>(KernelFunc, Kernel, {} /*Props*/, NumWorkItems,
WorkItemOffset);
}

/// Defines and invokes a SYCL kernel function for the specified range and
/// offsets.
///
/// \param Kernel is a SYCL kernel that is executed on a SYCL device
/// (except for the host device).
/// \param NDRange is a ND-range defining global and local sizes as
/// well as offset.
/// \param KernelFunc is a lambda that is used if device, queue is bound to,
/// 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_legacy<detail::WrapAs::parallel_for, KernelName, LambdaArgType,
Dims>(KernelFunc, Kernel, {} /*Props*/, NDRange);
}

/// Hierarchical kernel invocation method of a kernel.
///
/// This version of \c parallel_for_work_group takes two parameters
/// representing the same kernel. The first one - \c Kernel - is a
/// compiled form of the second one - \c kernelFunc, which is the source form
/// of the kernel. The same source kernel can be compiled multiple times
/// yielding multiple kernel class objects accessible via the \c program class
/// interface.
///
/// \param Kernel is a compiled SYCL kernel.
/// \param NumWorkGroups is a range describing the number of work-groups in
/// each dimension.
/// \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_legacy<detail::WrapAs::parallel_for_work_group, KernelName,
LambdaArgType, Dims,
/*SetNumWorkGroups*/ true>(KernelFunc, Kernel,
{} /*Props*/, NumWorkGroups);
}

/// Hierarchical kernel invocation method of a kernel.
///
/// This version of \c parallel_for_work_group takes two parameters
/// representing the same kernel. The first one - \c Kernel - is a
/// compiled form of the second one - \c kernelFunc, which is the source form
/// of the kernel. The same source kernel can be compiled multiple times
/// yielding multiple kernel class objects accessible via the \c program class
/// interface.
///
/// \param Kernel is a compiled SYCL kernel.
/// \param NumWorkGroups is a range describing the number of work-groups in
/// each dimension.
/// \param WorkGroupSize is a range describing the size of work-groups in
/// each dimension.
/// \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_legacy<detail::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>
__SYCL_DEPRECATED("To specify properties, use a launch configuration object "
Expand Down
9 changes: 0 additions & 9 deletions sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -869,15 +869,6 @@ void handler::extractArgsAndReqs() {
impl->MKernelData.extractArgsAndReqs(MKernel->isCreatedFromSource());
}

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// Calling methods of kernel_impl requires knowledge of class layout.
// As this is impossible in header, there's a function that calls necessary
// method inside the library and returns the result.
detail::ABINeutralKernelNameStrT handler::getKernelName() {
return MKernel->getName();
}
#endif

void handler::verifyUsedKernelBundleInternal(detail::string_view KernelName) {
detail::kernel_bundle_impl *UsedKernelBundleImplPtr =
getOrInsertHandlerKernelBundlePtr(/*Insert=*/false);
Expand Down
1 change: 0 additions & 1 deletion sycl/test/abi/sycl_symbols_linux.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3529,7 +3529,6 @@ _ZN4sycl3_V17handler11storeRawArgEPKvm
_ZN4sycl3_V17handler12addReductionERKSt10shared_ptrIKvE
_ZN4sycl3_V17handler12setArgHelperEiONS0_6streamE
_ZN4sycl3_V17handler12setArgHelperEiRNS0_6detail22work_group_memory_implE
_ZN4sycl3_V17handler13getKernelNameEv
_ZN4sycl3_V17handler13setKernelFuncEPv
_ZN4sycl3_V17handler15ext_oneapi_copyENS0_3ext6oneapi12experimental16image_mem_handleENS0_5rangeILi3EEERKNS4_16image_descriptorEPvS7_S7_S7_
_ZN4sycl3_V17handler15ext_oneapi_copyENS0_3ext6oneapi12experimental16image_mem_handleENS0_5rangeILi3EEERKNS4_16image_descriptorEPvS7_SA_mS7_
Expand Down
1 change: 0 additions & 1 deletion sycl/test/abi/sycl_symbols_windows.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4031,7 +4031,6 @@
?getElementSize@image_plain@detail@_V1@sycl@@IEBA_KXZ
?getEndTime@HostProfilingInfo@detail@_V1@sycl@@QEBA_KXZ
?getKernelBundle@handler@_V1@sycl@@AEBA?AV?$kernel_bundle@$0A@@23@XZ
?getKernelName@handler@_V1@sycl@@AEAA?AVstring@detail@23@XZ
?getMaxWorkGroups@handler@_V1@sycl@@AEAA?AV?$optional@V?$array@_K$02@std@@@std@@XZ
?getMaxWorkGroups_v2@handler@_V1@sycl@@AEAA?AV?$tuple@V?$array@_K$02@std@@_N@std@@XZ
?getMemoryObject@AccessorBaseHost@detail@_V1@sycl@@QEBAPEAXXZ
Expand Down