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

Author: sergey-semenov

sycl/include/sycl/handler.hpp

@@ -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);
 
@@ -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.
 
@@ -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 "

sycl/source/handler.cpp

@@ -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);

sycl/test/abi/sycl_symbols_linux.dump

@@ -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_

sycl/test/abi/sycl_symbols_windows.dump

@@ -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