[SYCL] Pass kernel bundle to SYCL 2020 interop kernel (#4877)

Interop kernels created using `make_kernel` does not have a reference to a kernel bundle, so calling `get_kernel_bundle` on these kernels will return an invalid kernel bundle. These changes pass the associated kernel bundle to the created interop kernel. This only applies to kernels created with `make_kernel`.

Author: steffenlarsen

sycl/source/backend.cpp

@@ -208,6 +208,7 @@ kernel make_kernel(const context &TargetContext,
                    backend Backend) {
   const auto &Plugin = getPlugin(Backend);
   const auto &ContextImpl = getSyclObjImpl(TargetContext);
+  const auto KernelBundleImpl = getSyclObjImpl(KernelBundle);
 
   // For Level-Zero expect exactly one device image in the bundle. This is
   // natural for interop kernel to get created out of a single native
@@ -218,7 +219,6 @@ kernel make_kernel(const context &TargetContext,
   //
   pi::PiProgram PiProgram = nullptr;
   if (Backend == backend::level_zero) {
-    auto KernelBundleImpl = getSyclObjImpl(KernelBundle);
     if (KernelBundleImpl->size() != 1)
       throw sycl::runtime_error{
           "make_kernel: kernel_bundle must have single program image",
@@ -241,7 +241,7 @@ kernel make_kernel(const context &TargetContext,
 
   // Construct the SYCL queue from PI queue.
   return detail::createSyclObjFromImpl<kernel>(
-      std::make_shared<kernel_impl>(PiKernel, ContextImpl));
+      std::make_shared<kernel_impl>(PiKernel, ContextImpl, KernelBundleImpl));
 }
 
 kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext,

sycl/source/detail/kernel_impl.cpp

@@ -17,10 +17,11 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
 
-kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context)
+kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context,
+                         KernelBundleImplPtr KernelBundleImpl)
     : kernel_impl(Kernel, Context,
                   std::make_shared<program_impl>(Context, Kernel),
-                  /*IsCreatedFromSource*/ true) {
+                  /*IsCreatedFromSource*/ true, KernelBundleImpl) {
   // This constructor is only called in the interoperability kernel constructor.
   // Let the runtime caller handle native kernel retaining in other cases if
   // it's needed.
@@ -35,11 +36,12 @@ kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context)
 }
 
 kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
-                         ProgramImplPtr ProgramImpl,
-                         bool IsCreatedFromSource)
+                         ProgramImplPtr ProgramImpl, bool IsCreatedFromSource,
+                         KernelBundleImplPtr KernelBundleImpl)
     : MKernel(Kernel), MContext(ContextImpl),
       MProgramImpl(std::move(ProgramImpl)),
-      MCreatedFromSource(IsCreatedFromSource) {
+      MCreatedFromSource(IsCreatedFromSource),
+      MKernelBundleImpl(std::move(KernelBundleImpl)) {
 
   RT::PiContext Context = nullptr;
   // Using the plugin from the passed ContextImpl
@@ -68,8 +70,7 @@ kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
   MIsInterop = MKernelBundleImpl->isInterop();
 }
 
-kernel_impl::kernel_impl(ContextImplPtr Context,
-                         ProgramImplPtr ProgramImpl)
+kernel_impl::kernel_impl(ContextImplPtr Context, ProgramImplPtr ProgramImpl)
     : MContext(Context), MProgramImpl(std::move(ProgramImpl)) {}
 
 kernel_impl::~kernel_impl() {

sycl/source/detail/kernel_impl.hpp

@@ -41,7 +41,9 @@ class kernel_impl {
   ///
   /// \param Kernel is a valid PiKernel instance
   /// \param Context is a valid SYCL context
-  kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context);
+  /// \param KernelBundleImpl is a valid instance of kernel_bundle_impl
+  kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context,
+              KernelBundleImplPtr KernelBundleImpl);
 
   /// Constructs a SYCL kernel instance from a SYCL program and a PiKernel
   ///
@@ -55,15 +57,17 @@ class kernel_impl {
   /// \param ProgramImpl is a valid instance of program_impl
   /// \param IsCreatedFromSource is a flag that indicates whether program
   /// is created from source code
+  /// \param KernelBundleImpl is a valid instance of kernel_bundle_impl
   kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
-              ProgramImplPtr ProgramImpl, bool IsCreatedFromSource);
+              ProgramImplPtr ProgramImpl, bool IsCreatedFromSource,
+              KernelBundleImplPtr KernelBundleImpl);
 
   /// Constructs a SYCL kernel_impl instance from a SYCL device_image,
   /// kernel_bundle and / PiKernel.
   ///
   /// \param Kernel is a valid PiKernel instance
   /// \param ContextImpl is a valid SYCL context
-  /// \param ProgramImpl is a valid instance of kernel_bundle_impl
+  /// \param KernelBundleImpl is a valid instance of kernel_bundle_impl
   kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
               DeviceImageImplPtr DeviceImageImpl,
               KernelBundleImplPtr KernelBundleImpl);

sycl/source/detail/program_impl.cpp

@@ -343,9 +343,9 @@ kernel program_impl::get_kernel(std::string KernelName,
     return createSyclObjFromImpl<kernel>(
         std::make_shared<kernel_impl>(MContext, PtrToSelf));
   }
-  return createSyclObjFromImpl<kernel>(std::make_shared<kernel_impl>(
-      get_pi_kernel(KernelName), MContext, PtrToSelf,
-      /*IsCreatedFromSource*/ IsCreatedFromSource));
+  return createSyclObjFromImpl<kernel>(
+      std::make_shared<kernel_impl>(get_pi_kernel(KernelName), MContext,
+                                    PtrToSelf, IsCreatedFromSource, nullptr));
 }
 
 std::vector<std::vector<char>> program_impl::get_binaries() const {

sycl/source/kernel.cpp

@@ -20,7 +20,7 @@ namespace sycl {
 kernel::kernel(cl_kernel ClKernel, const context &SyclContext)
     : impl(std::make_shared<detail::kernel_impl>(
           detail::pi::cast<detail::RT::PiKernel>(ClKernel),
-          detail::getSyclObjImpl(SyclContext))) {}
+          detail::getSyclObjImpl(SyclContext), nullptr)) {}
 
 cl_kernel kernel::get() const { return impl->get(); }