[SYCL][ABI-Break][PI] Remove deprecated piEnqueueNativeKernel (#10008)

This patch removes the deprecated entry-point `piEnqueueNativeKernel`.
This change is being made in preparation for porting the OpenCL plugin
to Unified Runtime which does not intend to support enqueuing native
kernels since it has already being deprecated.

Author: kbenzie

sycl/include/sycl/detail/cg.hpp

@@ -61,7 +61,6 @@ class CG {
     BarrierWaitlist = 6,
     Fill = 7,
     UpdateHost = 8,
-    RunOnHostIntel = 9,
     CopyUSM = 10,
     FillUSM = 11,
     PrefetchUSM = 12,
@@ -190,8 +189,7 @@ class CGExecKernel : public CG {
         MKernelName(std::move(KernelName)), MStreams(std::move(Streams)),
         MAuxiliaryResources(std::move(AuxiliaryResources)),
         MKernelCacheConfig(std::move(KernelCacheConfig)) {
-    assert((getType() == RunOnHostIntel || getType() == Kernel) &&
-           "Wrong type of exec kernel CG.");
+    assert(getType() == Kernel && "Wrong type of exec kernel CG.");
   }
 
   CGExecKernel(const CGExecKernel &CGExec) = default;

sycl/include/sycl/detail/pi.def

@@ -106,7 +106,6 @@ _PI_API(piSamplerRetain)
 _PI_API(piSamplerRelease)
 // Queue commands
 _PI_API(piEnqueueKernelLaunch)
-_PI_API(piEnqueueNativeKernel)
 _PI_API(piEnqueueEventsWait)
 _PI_API(piEnqueueEventsWaitWithBarrier)
 _PI_API(piEnqueueMemBufferRead)

sycl/include/sycl/detail/pi.h

@@ -1636,12 +1636,6 @@ __SYCL_EXPORT pi_result piEnqueueKernelLaunch(
     const size_t *local_work_size, pi_uint32 num_events_in_wait_list,
     const pi_event *event_wait_list, pi_event *event);
 
-__SYCL_EXPORT pi_result piEnqueueNativeKernel(
-    pi_queue queue, void (*user_func)(void *), void *args, size_t cb_args,
-    pi_uint32 num_mem_objects, const pi_mem *mem_list,
-    const void **args_mem_loc, pi_uint32 num_events_in_wait_list,
-    const pi_event *event_wait_list, pi_event *event);
-
 __SYCL_EXPORT pi_result piEnqueueEventsWait(pi_queue command_queue,
                                             pi_uint32 num_events_in_wait_list,
                                             const pi_event *event_wait_list,

sycl/include/sycl/handler.hpp

@@ -1696,24 +1696,6 @@ class __SYCL_EXPORT handler {
         std::move(KernelFunc));
   }
 
-  /// Defines and invokes a SYCL kernel on host device.
-  ///
-  /// \param Func is a SYCL kernel function defined by lambda function or a
-  /// named function object type.
-  template <typename FuncT>
-  __SYCL_DEPRECATED(
-      "run_on_host_intel() is deprecated, use host_task() instead")
-  void run_on_host_intel(FuncT Func) {
-    throwIfActionIsCreated();
-    // No need to check if range is out of INT_MAX limits as it's compile-time
-    // known constant
-    MNDRDesc.set(range<1>{1});
-
-    MArgs = std::move(MAssociatedAccesors);
-    MHostKernel.reset(new detail::HostKernel<FuncT, void, 1>(std::move(Func)));
-    setType(detail::CG::RunOnHostIntel);
-  }
-
   /// Enqueues a command to the SYCL runtime to invoke \p Func once.
   template <typename FuncT>
   std::enable_if_t<detail::check_fn_signature<std::remove_reference_t<FuncT>,

sycl/plugins/cuda/pi_cuda.cpp

@@ -139,7 +139,6 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piSamplerRelease, pi2ur::piSamplerRelease)
   // Queue commands
   _PI_CL(piEnqueueKernelLaunch, pi2ur::piEnqueueKernelLaunch)
-  _PI_CL(piEnqueueNativeKernel, pi2ur::piEnqueueNativeKernel)
   _PI_CL(piEnqueueEventsWait, pi2ur::piEnqueueEventsWait)
   _PI_CL(piEnqueueEventsWaitWithBarrier, pi2ur::piEnqueueEventsWaitWithBarrier)
   _PI_CL(piEnqueueMemBufferRead, pi2ur::piEnqueueMemBufferRead)

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -1919,12 +1919,6 @@ pi_result piextKernelGetNativeHandle(pi_kernel, pi_native_handle *) {
   DIE_NO_IMPLEMENTATION;
 }
 
-pi_result piEnqueueNativeKernel(pi_queue, void (*)(void *), void *, size_t,
-                                pi_uint32, const pi_mem *, const void **,
-                                pi_uint32, const pi_event *, pi_event *) {
-  DIE_NO_IMPLEMENTATION;
-}
-
 pi_result piextGetDeviceFunctionPointer(pi_device, pi_program, const char *,
                                         pi_uint64 *) {
   DIE_NO_IMPLEMENTATION;

sycl/plugins/hip/pi_hip.cpp

@@ -3178,30 +3178,6 @@ pi_result hip_piEnqueueKernelLaunch(
   return retError;
 }
 
-/// \TODO Not implemented
-pi_result
-hip_piEnqueueNativeKernel(pi_queue queue, void (*user_func)(void *), void *args,
-                          size_t cb_args, pi_uint32 num_mem_objects,
-                          const pi_mem *mem_list, const void **args_mem_loc,
-                          pi_uint32 num_events_in_wait_list,
-                          const pi_event *event_wait_list, pi_event *event) {
-  (void)queue;
-  (void)user_func;
-  (void)args;
-  (void)cb_args;
-  (void)num_mem_objects;
-  (void)mem_list;
-  (void)args_mem_loc;
-  (void)num_events_in_wait_list;
-  (void)event_wait_list;
-  (void)event;
-
-  sycl::detail::pi::die("Not implemented in HIP backend");
-  return {};
-}
-
-/// \TODO Not implemented
-
 pi_result hip_piMemImageCreate(pi_context context, pi_mem_flags flags,
                                const pi_image_format *image_format,
                                const pi_image_desc *image_desc, void *host_ptr,
@@ -6001,7 +5977,6 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piSamplerRelease, hip_piSamplerRelease)
   // Queue commands
   _PI_CL(piEnqueueKernelLaunch, hip_piEnqueueKernelLaunch)
-  _PI_CL(piEnqueueNativeKernel, hip_piEnqueueNativeKernel)
   _PI_CL(piEnqueueEventsWait, hip_piEnqueueEventsWait)
   _PI_CL(piEnqueueEventsWaitWithBarrier, hip_piEnqueueEventsWaitWithBarrier)
   _PI_CL(piEnqueueMemBufferRead, hip_piEnqueueMemBufferRead)

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -723,18 +723,6 @@ pi_result piMemBufferPartition(pi_mem Buffer, pi_mem_flags Flags,
                                      BufferCreateInfo, RetMem);
 }
 
-pi_result piEnqueueNativeKernel(pi_queue Queue, void (*UserFunc)(void *),
-                                void *Args, size_t CbArgs,
-                                pi_uint32 NumMemObjects, const pi_mem *MemList,
-                                const void **ArgsMemLoc,
-                                pi_uint32 NumEventsInWaitList,
-                                const pi_event *EventWaitList,
-                                pi_event *Event) {
-  return pi2ur::piEnqueueNativeKernel(
-      Queue, UserFunc, Args, CbArgs, NumMemObjects, MemList, ArgsMemLoc,
-      NumEventsInWaitList, EventWaitList, Event);
-}
-
 // TODO: Check if the function_pointer_ret type can be converted to void**.
 pi_result piextGetDeviceFunctionPointer(pi_device Device, pi_program Program,
                                         const char *FunctionName,

sycl/plugins/opencl/pi_opencl.cpp

@@ -2680,7 +2680,6 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piSamplerRelease, clReleaseSampler)
   // Queue commands
   _PI_CL(piEnqueueKernelLaunch, clEnqueueNDRangeKernel)
-  _PI_CL(piEnqueueNativeKernel, clEnqueueNativeKernel)
   _PI_CL(piEnqueueEventsWait, clEnqueueMarkerWithWaitList)
   _PI_CL(piEnqueueEventsWaitWithBarrier, clEnqueueBarrierWithWaitList)
   _PI_CL(piEnqueueMemBufferRead, clEnqueueReadBuffer)

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -1964,28 +1964,6 @@ piEnqueueMemImageFill(pi_queue Queue, pi_mem Image, const void *FillColor,
   return PI_SUCCESS;
 }
 
-inline pi_result
-piEnqueueNativeKernel(pi_queue Queue, void (*UserFunc)(void *), void *Args,
-                      size_t CbArgs, pi_uint32 NumMemObjects,
-                      const pi_mem *MemList, const void **ArgsMemLoc,
-                      pi_uint32 NumEventsInWaitList,
-                      const pi_event *EventsWaitList, pi_event *Event) {
-  std::ignore = UserFunc;
-  std::ignore = Args;
-  std::ignore = CbArgs;
-  std::ignore = NumMemObjects;
-  std::ignore = MemList;
-  std::ignore = ArgsMemLoc;
-  std::ignore = NumEventsInWaitList;
-  std::ignore = EventsWaitList;
-  std::ignore = Event;
-
-  PI_ASSERT(Queue, PI_ERROR_INVALID_QUEUE);
-
-  die("piEnqueueNativeKernel: not implemented");
-  return PI_SUCCESS;
-}
-
 inline pi_result piextGetDeviceFunctionPointer(pi_device Device,
                                                pi_program Program,
                                                const char *FunctionName,

sycl/plugins/unified_runtime/pi_unified_runtime.cpp

@@ -722,16 +722,6 @@ __SYCL_EXPORT pi_result piMemRelease(pi_mem Mem) {
   return pi2ur::piMemRelease(Mem);
 }
 
-__SYCL_EXPORT pi_result piEnqueueNativeKernel(
-    pi_queue Queue, void (*UserFunc)(void *), void *Args, size_t CbArgs,
-    pi_uint32 NumMemObjects, const pi_mem *MemList, const void **ArgsMemLoc,
-    pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList,
-    pi_event *Event) {
-  return pi2ur::piEnqueueNativeKernel(
-      Queue, UserFunc, Args, CbArgs, NumMemObjects, MemList, ArgsMemLoc,
-      NumEventsInWaitList, EventWaitList, Event);
-}
-
 __SYCL_EXPORT pi_result piextGetDeviceFunctionPointer(
     pi_device Device, pi_program Program, const char *FunctionName,
     pi_uint64 *FunctionPointerRet) {
@@ -1252,7 +1242,6 @@ __SYCL_EXPORT pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_API(piEnqueueMemBufferRead)
   _PI_API(piEnqueueEventsWaitWithBarrier)
   _PI_API(piEnqueueEventsWait)
-  _PI_API(piEnqueueNativeKernel)
   _PI_API(piEnqueueMemImageFill)
 
   _PI_API(piEventSetCallback)

sycl/source/detail/graph_impl.hpp

@@ -121,7 +121,6 @@ class node_impl {
   std::unique_ptr<sycl::detail::CG> getCGCopy() const {
     switch (MCGType) {
     case sycl::detail::CG::Kernel:
-    case sycl::detail::CG::RunOnHostIntel:
       return createCGCopy<sycl::detail::CGExecKernel>();
     case sycl::detail::CG::CopyAccToPtr:
     case sycl::detail::CG::CopyPtrToAcc:

sycl/source/detail/scheduler/commands.cpp

@@ -2320,8 +2320,6 @@ static pi_result SetKernelParamsAndLaunch(
 }
 
 // The function initialize accessors and calls lambda.
-// The function is used as argument to piEnqueueNativeKernel which requires
-// that the passed function takes one void* argument.
 void DispatchNativeKernel(void *Blob) {
   void **CastedBlob = (void **)Blob;
 
@@ -2596,93 +2594,6 @@ pi_int32 ExecCGCommand::enqueueImp() {
 
     return PI_SUCCESS;
   }
-  case CG::CGTYPE::RunOnHostIntel: {
-    CGExecKernel *HostTask = (CGExecKernel *)MCommandGroup.get();
-
-    // piEnqueueNativeKernel takes arguments blob which is passes to user
-    // function.
-    // Need the following items to restore context in the host task.
-    // Make a copy on heap to "dettach" from the command group as it can be
-    // released before the host task completes.
-    std::vector<void *> ArgsBlob(HostTask->MArgs.size() + 3);
-
-    std::vector<Requirement *> *CopyReqs =
-        new std::vector<Requirement *>(HostTask->getRequirements());
-
-    // Create a shared_ptr on the heap so that the reference count is
-    // incremented until the DispatchNativeKernel() callback is run, which
-    // will free the heap shared_ptr and decrement the reference count. This
-    // prevents errors when the HostTask command-group is deleted before
-    // DispatchNativeKernel() can be run.
-    std::shared_ptr<HostKernelBase> *CopyHostKernel =
-        new std::shared_ptr<HostKernelBase>(HostTask->MHostKernel);
-
-    NDRDescT *CopyNDRDesc = new NDRDescT(HostTask->MNDRDesc);
-
-    ArgsBlob[0] = (void *)CopyReqs;
-    ArgsBlob[1] = (void *)CopyHostKernel;
-    ArgsBlob[2] = (void *)CopyNDRDesc;
-
-    void **NextArg = ArgsBlob.data() + 3;
-
-    if (MQueue->is_host()) {
-      for (ArgDesc &Arg : HostTask->MArgs) {
-        assert(Arg.MType == kernel_param_kind_t::kind_accessor);
-
-        Requirement *Req = (Requirement *)(Arg.MPtr);
-        AllocaCommandBase *AllocaCmd = getAllocaForReq(Req);
-
-        *NextArg = AllocaCmd->getMemAllocation();
-        NextArg++;
-      }
-
-      if (!RawEvents.empty()) {
-        // Assuming that the events are for devices to the same Plugin.
-        const PluginPtr &Plugin = EventImpls[0]->getPlugin();
-        Plugin->call<PiApiKind::piEventsWait>(RawEvents.size(), &RawEvents[0]);
-      }
-      DispatchNativeKernel((void *)ArgsBlob.data());
-
-      return PI_SUCCESS;
-    }
-
-    std::vector<pi_mem> Buffers;
-    // piEnqueueNativeKernel requires additional array of pointers to args
-    // blob, values that pointers point to are replaced with actual pointers
-    // to the memory before execution of user function.
-    std::vector<void *> MemLocs;
-
-    for (ArgDesc &Arg : HostTask->MArgs) {
-      assert(Arg.MType == kernel_param_kind_t::kind_accessor);
-
-      Requirement *Req = (Requirement *)(Arg.MPtr);
-      AllocaCommandBase *AllocaCmd = getAllocaForReq(Req);
-      pi_mem MemArg = (pi_mem)AllocaCmd->getMemAllocation();
-
-      Buffers.push_back(MemArg);
-      MemLocs.push_back(NextArg);
-      NextArg++;
-    }
-    const PluginPtr &Plugin = MQueue->getPlugin();
-    pi_result Error = Plugin->call_nocheck<PiApiKind::piEnqueueNativeKernel>(
-        MQueue->getHandleRef(), DispatchNativeKernel, (void *)ArgsBlob.data(),
-        ArgsBlob.size() * sizeof(ArgsBlob[0]), Buffers.size(), Buffers.data(),
-        const_cast<const void **>(MemLocs.data()), RawEvents.size(),
-        RawEvents.empty() ? nullptr : RawEvents.data(), Event);
-
-    switch (Error) {
-    case PI_ERROR_INVALID_OPERATION:
-      throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
-                            "Device doesn't support run_on_host_intel tasks. " +
-                                detail::codeToString(Error));
-    case PI_SUCCESS:
-      return Error;
-    default:
-      throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
-                            "Enqueueing run_on_host_intel task has failed. " +
-                                detail::codeToString(Error));
-    }
-  }
   case CG::CGTYPE::Kernel: {
     CGExecKernel *ExecKernel = (CGExecKernel *)MCommandGroup.get();
 

sycl/source/handler.cpp

@@ -273,8 +273,7 @@ event handler::finalize() {
 
   std::unique_ptr<detail::CG> CommandGroup;
   switch (type) {
-  case detail::CG::Kernel:
-  case detail::CG::RunOnHostIntel: {
+  case detail::CG::Kernel: {
     // Copy kernel name here instead of move so that it's available after
     // running of this method by reductions implementation. This allows for
     // assert feature to check if kernel uses assertions