[SYCL] Implement DeviceKernelInfo for interop kernels

sycl/include/sycl/handler.hpp

@@ -501,14 +501,7 @@ class __SYCL_EXPORT handler {
     setType(detail::CGType::Kernel);
   }
 
-  void setDeviceKernelInfo(kernel &&Kernel) {
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
-    MKernelName = getKernelName();
-    setType(detail::CGType::Kernel);
-
-    // If any extra actions are added here make sure that logic around
-    // `lambdaAndKernelHaveEqualName` calls can handle that.
-  }
+  void setDeviceKernelInfo(kernel &&Kernel);
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   // TODO: Those functions are not used anymore, remove it in the next

sycl/source/detail/kernel_impl.cpp

@@ -26,7 +26,9 @@ kernel_impl::kernel_impl(Managed<ur_kernel_handle_t> &&Kernel,
       MCreatedFromSource(true),
       MKernelBundleImpl(KernelBundleImpl ? KernelBundleImpl->shared_from_this()
                                          : nullptr),
-      MIsInterop(true), MKernelArgMaskPtr{ArgMask} {
+      MIsInterop(true), MKernelArgMaskPtr{ArgMask},
+      MInteropDeviceKernelInfoHolder(CompileTimeKernelInfoTy{getName()}),
+      MDeviceKernelInfo(MInteropDeviceKernelInfoHolder) {
   ur_context_handle_t UrContext = nullptr;
   // Using the adapter from the passed ContextImpl
   getAdapter().call<UrApiKind::urKernelGetInfo>(
@@ -52,7 +54,15 @@ kernel_impl::kernel_impl(Managed<ur_kernel_handle_t> &&Kernel,
       MDeviceImageImpl(std::move(DeviceImageImpl)),
       MKernelBundleImpl(KernelBundleImpl.shared_from_this()),
       MIsInterop(MDeviceImageImpl->getOriginMask() & ImageOriginInterop),
-      MKernelArgMaskPtr{ArgMask}, MCacheMutex{CacheMutex} {
+      MKernelArgMaskPtr{ArgMask}, MCacheMutex{CacheMutex},
+      MInteropDeviceKernelInfoHolder(MIsInterop
+                                         ? CompileTimeKernelInfoTy{getName()}
+                                         : CompileTimeKernelInfoTy{}),
+      MDeviceKernelInfo(
+          MIsInterop
+              ? MInteropDeviceKernelInfoHolder
+              : ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
+                    KernelNameStrT(getName()))) {
   // Enable USM indirect access for interop and non-sycl-jit source kernels.
   // sycl-jit kernels will enable this if needed through the regular kernel
   // path.

sycl/source/detail/kernel_impl.hpp

@@ -238,6 +238,8 @@ class kernel_impl {
   std::mutex *getCacheMutex() const { return MCacheMutex; }
   std::string_view getName() const;
 
+  DeviceKernelInfo &getDeviceKernelInfo() const { return MDeviceKernelInfo; }
+
 private:
   Managed<ur_kernel_handle_t> MKernel;
   const std::shared_ptr<context_impl> MContext;
@@ -251,6 +253,12 @@ class kernel_impl {
   std::mutex *MCacheMutex = nullptr;
   mutable std::string MName;
 
+  // For the interop kernel we create the DeviceKernelInfo
+  // as part of the kernel_impl.
+  // For regular kernel we get DeviceKernelInfo from the ProgramManager.
+  DeviceKernelInfo MInteropDeviceKernelInfoHolder;
+  DeviceKernelInfo &MDeviceKernelInfo;
+
   bool isBuiltInKernel(device_impl &Device) const;
   void checkIfValidForNumArgsInfoQuery() const;
 

sycl/source/handler.cpp

@@ -2535,6 +2535,16 @@ bool handler::HasAssociatedAccessor(detail::AccessorImplHost *Req,
 void handler::setType(sycl::detail::CGType Type) { impl->MCGType = Type; }
 sycl::detail::CGType handler::getType() const { return impl->MCGType; }
 
+void handler::setDeviceKernelInfo(kernel &&Kernel) {
+  MKernel = detail::getSyclObjImpl(std::move(Kernel));
+  MKernelName = getKernelName();
+  setDeviceKernelInfoPtr(&MKernel->getDeviceKernelInfo());
+  setType(detail::CGType::Kernel);
+
+  // If any extra actions are added here make sure that logic around
+  // `lambdaAndKernelHaveEqualName` calls can handle that.
+}
+
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 void handler::setNDRangeDescriptorPadded(sycl::range<3> N,
                                          bool SetNumWorkGroups, int Dims) {

sycl/test/abi/sycl_symbols_linux.dump

@@ -3592,6 +3592,7 @@ _ZN4sycl3_V17handler21setKernelWorkGroupMemEm
 _ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE
 _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm
 _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm
+_ZN4sycl3_V17handler19setDeviceKernelInfoEONS0_6kernelE
 _ZN4sycl3_V17handler22setDeviceKernelInfoPtrEPNS0_6detail16DeviceKernelInfoE
 _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE
 _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE