[SYCL][L0] Add ownership control for LeveL-Zero kernel_bundle interop. (#4576)

Signed-off-by: Sergey V Maslov <sergey.v.maslov@intel.com>

Author: smaslov-intel

sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md

@@ -101,7 +101,8 @@ backend_input_t<backend::ext_oneapi_level_zero,
 struct {
   ze_context_handle_t NativeHandle;
   std::vector<device> DeviceList;
-  ext::oneapi::level_zero::ownership Ownership;
+  ext::oneapi::level_zero::ownership Ownership{
+      ext::oneapi::level_zero::ownership::transfer};
 }
 ```
 </td>
@@ -113,7 +114,8 @@ struct {
 ``` C++
 struct {
   ze_command_queue_handle_t NativeHandle;
-  ext::oneapi::level_zero::ownership Ownership;
+  ext::oneapi::level_zero::ownership Ownership{
+      ext::oneapi::level_zero::ownership::transfer};
 }
 ```
 </td>
@@ -125,7 +127,8 @@ struct {
 ``` C++
 struct {
   ze_event_handle_t NativeHandle;
-  ext::oneapi::level_zero::ownership Ownership;
+  ext::oneapi::level_zero::ownership Ownership{
+      ext::oneapi::level_zero::ownership::transfer};
 }
 ```
 </td>
@@ -137,7 +140,16 @@ struct {
 std::vector<ze_module_handle_t>
 ```
 </td>
-<td><pre>ze_module_handle_t</pre></td>
+<td>
+
+``` C++
+struct {
+  ze_module_handle_t NativeHandle;
+  ext::oneapi::level_zero::ownership Ownership{
+      ext::oneapi::level_zero::ownership::transfer};
+}
+```
+</td>
 </tr>
 </table>
 
@@ -226,7 +238,20 @@ make_kernel_bundle<backend::ext_oneapi_level_zero,
     const context &Context)
 ```
 </td>
-<td>Constructs a SYCL kernel_bundle instance from a Level-Zero <code>ze_module_handle_t</code>. The <code>Context</code> argument must be a valid SYCL context encapsulating a Level-Zero context. The Level-Zero module must be fully linked (i.e. not require further linking through <a href="https://spec.oneapi.com/level-zero/latest/core/api.html?highlight=zemoduledynamiclink#_CPPv419zeModuleDynamicLink8uint32_tP18ze_module_handle_tP28ze_module_build_log_handle_t"><code>zeModuleDynamicLink</code></a>), and thus the SYCL kernel_bundle is created in the "executable" state.</td>
+<td>Constructs a SYCL kernel_bundle instance from a Level-Zero
+<code>ze_module_handle_t</code>. The <code>Context</code> argument must be a
+valid SYCL context encapsulating a Level-Zero context, and the Level-Zero
+module must be created on the same context. The Level-Zero module must be
+fully linked (i.e. not require further linking through <a href="https://spec.oneapi.com/level-zero/latest/core/api.html?highlight=zemoduledynamiclink#_CPPv419zeModuleDynamicLink8uint32_tP18ze_module_handle_tP28ze_module_build_log_handle_t">
+<code>zeModuleDynamicLink</code></a>), and thus the SYCL kernel_bundle is
+created in the "executable" state. The <code>Ownership</code> input structure
+member specifies if the SYCL runtime should take ownership of the passed
+native handle. The default behavior is to transfer the ownership to the SYCL
+runtime. See section 4.4 for details. If the behavior is "transfer" then the
+runtime is going to destroy the input Level-Zero module, and hence the
+application must not to have any outstanding <code>ze_kernel_handle_t</code>
+handles to the underlying <code>ze_module_handle_t</code> by the time this
+interoperability <code>kernel_bundle</code> destructor is called.</td>
 </tr>
 </table>
 

sycl/include/CL/sycl/backend.hpp

@@ -108,9 +108,13 @@ __SYCL_EXPORT event make_event(pi_native_handle NativeHandle,
                                backend Backend);
 __SYCL_EXPORT kernel make_kernel(pi_native_handle NativeHandle,
                                  const context &TargetContext, backend Backend);
+// TODO: Unused. Remove when allowed.
 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
 make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
                    bundle_state State, backend Backend);
+__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
+make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
+                   bool KeepOwnership, bundle_state State, backend Backend);
 } // namespace detail
 
 template <backend Backend>
@@ -221,7 +225,7 @@ make_kernel_bundle(const typename backend_traits<Backend>::template input_type<
   std::shared_ptr<detail::kernel_bundle_impl> KBImpl =
       detail::make_kernel_bundle(
           detail::pi::cast<pi_native_handle>(BackendObject), TargetContext,
-          State, Backend);
+          false, State, Backend);
   return detail::createSyclObjFromImpl<kernel_bundle<State>>(KBImpl);
 }
 } // namespace sycl

sycl/include/CL/sycl/detail/pi.h

@@ -40,6 +40,7 @@
 // changes the API version from 3.5 to 4.6.
 // 5.7 Added new context and ownership arguments to
 //   piextEventCreateWithNativeHandle
+// 6.8 Added new ownership argument to piextProgramCreateWithNativeHandle.
 //
 #include "CL/cl.h"
 #define _PI_H_VERSION_MAJOR 5
@@ -1033,8 +1034,8 @@ piextContextGetNativeHandle(pi_context context, pi_native_handle *nativeHandle);
 /// \param devices is the list of devices in the context. Parameter is ignored
 ///        if devices can be queried from the context native handle for a
 ///        backend.
-/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
-///        the native handle, if it can.
+/// \param pluginOwnsNativeHandle Indicates whether the created PI object
+///        should take ownership of the native handle.
 /// \param context is the PI context created from the native handle.
 /// \return PI_SUCCESS if successfully created pi_context from the handle.
 ///         PI_OUT_OF_HOST_MEMORY if can't allocate memory for the pi_context
@@ -1043,7 +1044,7 @@ piextContextGetNativeHandle(pi_context context, pi_native_handle *nativeHandle);
 ///         native handle. PI_UNKNOWN_ERROR in case of another error.
 __SYCL_EXPORT pi_result piextContextCreateWithNativeHandle(
     pi_native_handle nativeHandle, pi_uint32 numDevices,
-    const pi_device *devices, bool ownNativeHandle, pi_context *context);
+    const pi_device *devices, bool pluginOwnsNativeHandle, pi_context *context);
 
 //
 // Queue
@@ -1077,11 +1078,11 @@ piextQueueGetNativeHandle(pi_queue queue, pi_native_handle *nativeHandle);
 /// \param nativeHandle is the native handle to create PI queue from.
 /// \param context is the PI context of the queue.
 /// \param queue is the PI queue created from the native handle.
-/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
-///        the native handle, if it can.
+/// \param pluginOwnsNativeHandle Indicates whether the created PI object
+///        should take ownership of the native handle.
 __SYCL_EXPORT pi_result piextQueueCreateWithNativeHandle(
     pi_native_handle nativeHandle, pi_context context, pi_queue *queue,
-    bool ownNativeHandle);
+    bool pluginOwnsNativeHandle);
 
 //
 // Memory
@@ -1219,9 +1220,12 @@ piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle);
 ///
 /// \param nativeHandle is the native handle to create PI program from.
 /// \param context is the PI context of the program.
+/// \param pluginOwnsNativeHandle Indicates whether the created PI object
+///        should take ownership of the native handle.
 /// \param program is the PI program created from the native handle.
 __SYCL_EXPORT pi_result piextProgramCreateWithNativeHandle(
-    pi_native_handle nativeHandle, pi_context context, pi_program *program);
+    pi_native_handle nativeHandle, pi_context context,
+    bool pluginOwnsNativeHandle, pi_program *program);
 
 //
 // Kernel
@@ -1315,12 +1319,12 @@ __SYCL_EXPORT pi_result piKernelSetExecInfo(pi_kernel kernel,
 ///
 /// \param nativeHandle is the native handle to create PI kernel from.
 /// \param context is the PI context of the kernel.
-/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
-///        the native handle, if it can.
+/// \param pluginOwnsNativeHandle Indicates whether the created PI object
+///        should take ownership of the native handle.
 /// \param kernel is the PI kernel created from the native handle.
 __SYCL_EXPORT pi_result piextKernelCreateWithNativeHandle(
-    pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,
-    pi_kernel *kernel);
+    pi_native_handle nativeHandle, pi_context context,
+    bool pluginOwnsNativeHandle, pi_kernel *kernel);
 
 /// Gets the native handle of a PI kernel object.
 ///
@@ -1373,8 +1377,8 @@ piextEventGetNativeHandle(pi_event event, pi_native_handle *nativeHandle);
 ///
 /// \param nativeHandle is the native handle to create PI event from.
 /// \param context is the corresponding PI context
-/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
-///        the native handle, if it can.
+/// \param pluginOwnsNativeHandle Indicates whether the created PI object
+///        should take ownership of the native handle.
 /// \param event is the PI event created from the native handle.
 __SYCL_EXPORT pi_result piextEventCreateWithNativeHandle(
     pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,

sycl/include/sycl/ext/oneapi/backend/level_zero.hpp

@@ -101,7 +101,11 @@ template <> struct BackendInput<backend::level_zero, event> {
 
 template <bundle_state State>
 struct BackendInput<backend::level_zero, kernel_bundle<State>> {
-  using type = ze_module_handle_t;
+  using type = struct {
+    ze_module_handle_t NativeHandle;
+    ext::oneapi::level_zero::ownership Ownership{
+        ext::oneapi::level_zero::ownership::transfer};
+  };
 };
 
 template <bundle_state State>
@@ -249,6 +253,24 @@ event make_event<backend::level_zero>(
       BackendObject.Ownership == ext::oneapi::level_zero::ownership::keep);
 }
 
+// Specialization of sycl::make_kernel_bundle for Level-Zero backend.
+template <>
+kernel_bundle<bundle_state::executable>
+make_kernel_bundle<backend::ext_oneapi_level_zero, bundle_state::executable>(
+    const backend_input_t<backend::ext_oneapi_level_zero,
+                          kernel_bundle<bundle_state::executable>>
+        &BackendObject,
+    const context &TargetContext) {
+  std::shared_ptr<detail::kernel_bundle_impl> KBImpl =
+      detail::make_kernel_bundle(
+          detail::pi::cast<pi_native_handle>(BackendObject.NativeHandle),
+          TargetContext,
+          BackendObject.Ownership == ext::oneapi::level_zero::ownership::keep,
+          bundle_state::executable, backend::ext_oneapi_level_zero);
+  return detail::createSyclObjFromImpl<kernel_bundle<bundle_state::executable>>(
+      KBImpl);
+}
+
 // TODO: remove this specialization when generic is changed to call
 // .GetNative() instead of .get_native() member of kernel_bundle.
 template <>

sycl/plugins/cuda/pi_cuda.cpp

@@ -3174,7 +3174,7 @@ pi_result cuda_piextProgramGetNativeHandle(pi_program program,
 ///
 /// \return TBD
 pi_result cuda_piextProgramCreateWithNativeHandle(pi_native_handle, pi_context,
-                                                  pi_program *) {
+                                                  bool, pi_program *) {
   cl::sycl::detail::pi::die(
       "Creation of PI program from native handle not implemented");
   return {};

sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp

@@ -1019,19 +1019,15 @@ pi_result piProgramGetBuildInfo(pi_program, pi_device, cl_program_build_info,
   DIE_NO_IMPLEMENTATION;
 }
 
-pi_result piProgramRetain(pi_program) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piProgramRetain(pi_program) { DIE_NO_IMPLEMENTATION; }
 
-pi_result piProgramRelease(pi_program) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piProgramRelease(pi_program) { DIE_NO_IMPLEMENTATION; }
 
 pi_result piextProgramGetNativeHandle(pi_program, pi_native_handle *) {
   DIE_NO_IMPLEMENTATION;
 }
 
-pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context,
+pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context, bool,
                                              pi_program *) {
   DIE_NO_IMPLEMENTATION;
 }
@@ -1068,17 +1064,11 @@ pi_result piKernelGetSubGroupInfo(pi_kernel, pi_device,
   DIE_NO_IMPLEMENTATION;
 }
 
-pi_result piKernelRetain(pi_kernel) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piKernelRetain(pi_kernel) { DIE_NO_IMPLEMENTATION; }
 
-pi_result piKernelRelease(pi_kernel) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piKernelRelease(pi_kernel) { DIE_NO_IMPLEMENTATION; }
 
-pi_result piEventCreate(pi_context, pi_event *) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piEventCreate(pi_context, pi_event *) { DIE_NO_IMPLEMENTATION; }
 
 pi_result piEventGetInfo(pi_event, pi_event_info, size_t, void *, size_t *) {
   DIE_NO_IMPLEMENTATION;
@@ -1117,9 +1107,7 @@ pi_result piEventSetCallback(pi_event, pi_int32,
   DIE_NO_IMPLEMENTATION;
 }
 
-pi_result piEventSetStatus(pi_event, pi_int32) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piEventSetStatus(pi_event, pi_int32) { DIE_NO_IMPLEMENTATION; }
 
 pi_result piEventRetain(pi_event Event) {
   if (Event == nullptr) {
@@ -1170,13 +1158,9 @@ pi_result piSamplerGetInfo(pi_sampler, pi_sampler_info, size_t, void *,
   DIE_NO_IMPLEMENTATION;
 }
 
-pi_result piSamplerRetain(pi_sampler) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piSamplerRetain(pi_sampler) { DIE_NO_IMPLEMENTATION; }
 
-pi_result piSamplerRelease(pi_sampler) {
-  DIE_NO_IMPLEMENTATION;
-}
+pi_result piSamplerRelease(pi_sampler) { DIE_NO_IMPLEMENTATION; }
 
 pi_result piEnqueueEventsWait(pi_queue, pi_uint32, const pi_event *,
                               pi_event *) {

sycl/plugins/hip/pi_hip.cpp

@@ -3018,11 +3018,14 @@ pi_result hip_piextProgramGetNativeHandle(pi_program program,
 ///
 /// \param[in] nativeHandle The native handle to create PI program object from.
 /// \param[in] context The PI context of the program.
+/// \param[in] ownNativeHandle tells if should assume the ownership of
+///            the native handle.
 /// \param[out] program Set to the PI program object created from native handle.
 ///
 /// \return TBD
 pi_result hip_piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
                                                  pi_context context,
+                                                 bool ownNativeHandle,
                                                  pi_program *program) {
   (void)nativeHandle;
   (void)context;

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -3657,8 +3657,9 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices,
           if (res != PI_SUCCESS) {
             return res;
           }
-          Input = new _pi_program(Input->Context, ZeModule, _pi_program::Object,
-                                  Input->HasImports);
+          Input =
+              new _pi_program(Input->Context, ZeModule, true /*own ZeModule*/,
+                              _pi_program::Object, Input->HasImports);
           Input->HasImportsAndIsLinked = true;
         }
       } else {
@@ -3913,6 +3914,7 @@ pi_result piextProgramGetNativeHandle(pi_program Program,
 
 pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
                                              pi_context Context,
+                                             bool ownNativeHandle,
                                              pi_program *Program) {
   PI_ASSERT(Program, PI_INVALID_PROGRAM);
   PI_ASSERT(NativeHandle, PI_INVALID_VALUE);
@@ -3925,7 +3927,8 @@ pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
   // executable (state Object).
 
   try {
-    *Program = new _pi_program(Context, ZeModule, _pi_program::Exe);
+    *Program =
+        new _pi_program(Context, ZeModule, ownNativeHandle, _pi_program::Exe);
   } catch (const std::bad_alloc &) {
     return PI_OUT_OF_HOST_MEMORY;
   } catch (...) {
@@ -3942,7 +3945,7 @@ _pi_program::~_pi_program() {
     ZE_CALL_NOCHECK(zeModuleBuildLogDestroy, (ZeBuildLog));
   }
 
-  if (ZeModule) {
+  if (ZeModule && OwnZeModule) {
     ZE_CALL_NOCHECK(zeModuleDestroy, (ZeModule));
   }
 }

sycl/plugins/level_zero/pi_level_zero.hpp

@@ -1068,23 +1068,24 @@ struct _pi_program : _pi_object {
   // Construct a program in IL or Native state.
   _pi_program(pi_context Context, const void *Input, size_t Length, state St)
       : State(St), Context(Context), Code(new uint8_t[Length]),
-        CodeLength(Length), ZeModule(nullptr), HasImports(false),
-        HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {
+        CodeLength(Length), ZeModule(nullptr), OwnZeModule{true},
+        HasImports(false), HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {
 
     std::memcpy(Code.get(), Input, Length);
   }
 
   // Construct a program in either Object or Exe state.
-  _pi_program(pi_context Context, ze_module_handle_t ZeModule, state St,
-              bool HasImports = false)
-      : State(St), Context(Context), ZeModule(ZeModule), HasImports(HasImports),
+  _pi_program(pi_context Context, ze_module_handle_t ZeModule, bool OwnZeModule,
+              state St, bool HasImports = false)
+      : State(St), Context(Context),
+        ZeModule(ZeModule), OwnZeModule{OwnZeModule}, HasImports(HasImports),
         HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {}
 
   // Construct a program in LinkedExe state.
   _pi_program(pi_context Context, std::vector<LinkedReleaser> &&Inputs,
               ze_module_build_log_handle_t ZeLog)
       : State(LinkedExe), Context(Context), ZeModule(nullptr),
-        HasImports(false), HasImportsAndIsLinked(false),
+        OwnZeModule(true), HasImports(false), HasImportsAndIsLinked(false),
         LinkedPrograms(std::move(Inputs)), ZeBuildLog(ZeLog) {}
 
   ~_pi_program();
@@ -1103,7 +1104,13 @@ struct _pi_program : _pi_object {
 
   // Used for programs in Object or Exe state.
   ze_module_handle_t ZeModule; // Level Zero module handle.
-  bool HasImports;             // Tells if module imports any symbols.
+
+  // Indicates if we own the ZeModule or it came from interop that
+  // asked to not transfer the ownership to SYCL RT.
+  bool OwnZeModule;
+
+  // Tells if module imports any symbols.
+  bool HasImports;
 
   // Used for programs in Object state.  Tells if this module imports any
   // symbols AND it is linked into some other program that has state LinkedExe.