diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index a0fcc39bfbb1..44b46d496091 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/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
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index cb78b8f50e1c..dfed6881a3ef 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -16,6 +16,11 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
+static CompileTimeKernelInfoTy
+createCompileTimeKernelInfo(std::string_view KernelName = {}) {
+  return CompileTimeKernelInfoTy{KernelName};
+}
+
 kernel_impl::kernel_impl(Managed<ur_kernel_handle_t> &&Kernel,
                          context_impl &Context,
                          kernel_bundle_impl *KernelBundleImpl,
@@ -26,7 +31,8 @@ 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},
+      MInteropDeviceKernelInfo(createCompileTimeKernelInfo(getName())) {
   ur_context_handle_t UrContext = nullptr;
   // Using the adapter from the passed ContextImpl
   getAdapter().call<UrApiKind::urKernelGetInfo>(
@@ -52,7 +58,10 @@ 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},
+      MInteropDeviceKernelInfo(MIsInterop
+                                   ? createCompileTimeKernelInfo(getName())
+                                   : createCompileTimeKernelInfo()) {
   // 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.
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index c6135f1fea9d..ae40e16fbe36 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -238,6 +238,13 @@ class kernel_impl {
   std::mutex *getCacheMutex() const { return MCacheMutex; }
   std::string_view getName() const;
 
+  DeviceKernelInfo &getDeviceKernelInfo() {
+    return MIsInterop
+               ? MInteropDeviceKernelInfo
+               : ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
+                     KernelNameStrT(getName()));
+  }
+
 private:
   Managed<ur_kernel_handle_t> MKernel;
   const std::shared_ptr<context_impl> MContext;
@@ -251,6 +258,10 @@ class kernel_impl {
   std::mutex *MCacheMutex = nullptr;
   mutable std::string MName;
 
+  // It is used for the interop kernels only.
+  // For regular kernel we get DeviceKernelInfo from the ProgramManager.
+  DeviceKernelInfo MInteropDeviceKernelInfo;
+
   bool isBuiltInKernel(device_impl &Device) const;
   void checkIfValidForNumArgsInfoQuery() const;
 
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index ddb09c5f7665..f11aba30f025 100644
--- a/sycl/source/handler.cpp
+++ b/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) {
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index e9f65ce66248..59c8bce5b885 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/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