From 1078abe61d811ec53347599515d446ab3b8205e8 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Mon, 24 Nov 2025 07:07:13 -0800 Subject: [PATCH 1/2] [SYCL][ABI-Break] Drop old enqueue functions from handler --- sycl/include/sycl/handler.hpp | 234 -------------------------- sycl/source/handler.cpp | 9 - sycl/test/abi/sycl_symbols_linux.dump | 1 - 3 files changed, 244 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index d8d46d2a27814..20e3dbf564ba2 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -528,24 +528,6 @@ class __SYCL_EXPORT handler { bool IsKernelCreatedFromSource, bool IsESIMD); #endif -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES - /// \return a string containing name of SYCL kernel. - detail::ABINeutralKernelNameStrT getKernelName(); - - template 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.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. #ifndef __INTEL_PREVIEW_BREAKING_CHANGES @@ -1371,64 +1353,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::name; - (void)Props; - (void)Kernel; - detail::KernelWrapper::wrap(KernelFunc); - - if constexpr (ext::oneapi::experimental::detail:: - HasKernelPropertiesGetMethod::value) { - SetKernelLaunchpropertiesIfNotEmpty(detail::extractKernelProperties( - KernelFunc.get(ext::oneapi::experimental::properties_tag{}))); - } - -#ifndef __SYCL_DEVICE_ONLY__ - constexpr auto Info = detail::CompileTimeKernelInfo; - 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(params...); - if constexpr (SetNumWorkGroups) { - setNDRangeDescriptor(std::move(params)..., - /*SetNumWorkGroups=*/true); - } else { - setNDRangeDescriptor(std::move(params)...); - } - - MKernel = detail::getSyclObjImpl(Kernel); - if (!lambdaAndKernelHaveEqualName()) { - throw sycl::exception( - make_error_code(errc::invalid), - "the kernel name must match the name of the lambda"); - } - StoreLambda(std::move(KernelFunc)); - SetKernelLaunchpropertiesIfNotEmpty( - detail::extractKernelProperties(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. @@ -1856,164 +1780,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 - __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::name; - (void)Kernel; - detail::KernelWrapperHelperFuncs::kernel_single_task(KernelFunc); -#ifndef __SYCL_DEVICE_ONLY__ - throwIfActionIsCreated(); - constexpr auto Info = detail::CompileTimeKernelInfo; - 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()) { - throw sycl::exception( - make_error_code(errc::invalid), - "the kernel name must match the name of the lambda"); - } - StoreLambda(std::move(KernelFunc)); - -#else - detail::CheckDeviceCopyable(); -#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 - __SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.") - void parallel_for(kernel Kernel, range 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>; - wrap_kernel_legacy(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 - __SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.") - void parallel_for(kernel Kernel, range NumWorkItems, - id WorkItemOffset, const KernelType &KernelFunc) { - using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy(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 - __SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.") - void parallel_for(kernel Kernel, nd_range NDRange, - const KernelType &KernelFunc) { - using LambdaArgType = - sycl::detail::lambda_arg_type>; - wrap_kernel_legacy(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 - __SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.") - void parallel_for_work_group(kernel Kernel, range NumWorkGroups, - const KernelType &KernelFunc) { - using LambdaArgType = - sycl::detail::lambda_arg_type>; - wrap_kernel_legacy(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 - __SYCL_DEPRECATED("This overload isn't part of SYCL2020 and will be removed.") - void parallel_for_work_group(kernel Kernel, range NumWorkGroups, - range WorkGroupSize, - const KernelType &KernelFunc) { - using LambdaArgType = - sycl::detail::lambda_arg_type>; - nd_range ExecRange = - nd_range(NumWorkGroups * WorkGroupSize, WorkGroupSize); - wrap_kernel_legacy(KernelFunc, Kernel, {} /*Props*/, - ExecRange); - } -#endif // __INTEL_PREVIEW_BREAKING_CHANGES - template __SYCL_DEPRECATED("To specify properties, use a launch configuration object " diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index fbcd88f1bd42a..85cc969b1aa2e 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -1050,15 +1050,6 @@ void handler::extractArgsAndReqsFromLambda( } #endif // __INTEL_PREVIEW_BREAKING_CHANGES -#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); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index ec52f5008c47a..c260d547b730b 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3553,7 +3553,6 @@ _ZN4sycl3_V17handler11storeRawArgEPKvm _ZN4sycl3_V17handler12addReductionERKSt10shared_ptrIKvE _ZN4sycl3_V17handler12setArgHelperEiONS0_6streamE _ZN4sycl3_V17handler12setArgHelperEiRNS0_6detail22work_group_memory_implE -_ZN4sycl3_V17handler13getKernelNameEv _ZN4sycl3_V17handler13setKernelFuncEPv _ZN4sycl3_V17handler13setKernelInfoEPviPFNS0_6detail19kernel_param_desc_tEiEbb _ZN4sycl3_V17handler14addAccessorReqESt10shared_ptrINS0_6detail16AccessorImplHostEE From 69c72e3069dd8f1e798f7ba29aefeb9194675a6b Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Tue, 25 Nov 2025 06:17:55 -0800 Subject: [PATCH 2/2] Update Windows ABI --- sycl/test/abi/sycl_symbols_windows.dump | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index cb6c4138f09eb..4d023330e4f4a 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4057,7 +4057,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