diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index b5f1f53e32d7f..42f26547a58df 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -1474,11 +1474,10 @@ void exec_graph_impl::populateURKernelUpdateStructs( ur_kernel_handle_t UrKernel = nullptr; auto Kernel = ExecCG.MSyclKernel; auto KernelBundleImplPtr = ExecCG.MKernelBundle; - std::shared_ptr SyclKernelImpl = nullptr; const sycl::detail::KernelArgMask *EliminatedArgMask = nullptr; if (auto SyclKernelImpl = KernelBundleImplPtr - ? KernelBundleImplPtr->tryGetKernel( + ? KernelBundleImplPtr->tryGetOfflineKernel( ExecCG.MKernelName, KernelBundleImplPtr) : std::shared_ptr{nullptr}) { UrKernel = SyclKernelImpl->getHandleRef(); diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 12ac9a47b7c07..183215c5aac0a 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -537,7 +537,7 @@ class kernel_bundle_impl { get_kernel(const kernel_id &KernelID, const std::shared_ptr &Self) const { if (std::shared_ptr KernelImpl = - tryGetOfflineKernel(KernelID, Self)) + tryGetOfflineKernelImpl(KernelID, Self)) return detail::createSyclObjFromImpl(std::move(KernelImpl)); throw sycl::exception(make_error_code(errc::invalid), "The kernel bundle does not contain the kernel " @@ -692,7 +692,7 @@ class kernel_bundle_impl { }); } - std::shared_ptr tryGetOfflineKernel( + std::shared_ptr tryGetOfflineKernelImpl( const kernel_id &KernelID, const std::shared_ptr &Self) const { using ImageImpl = std::shared_ptr; @@ -758,6 +758,17 @@ class kernel_bundle_impl { SelectedImage->get_ur_program_ref(), CacheMutex); } + std::shared_ptr + tryGetOfflineKernel(detail::KernelNameStrRefT Name, + const std::shared_ptr &Self) const { + // Fall back to regular offline compiled kernel_bundle look-up. + if (std::optional MaybeKernelID = + sycl::detail::ProgramManager::getInstance().tryGetSYCLKernelID( + Name)) + return tryGetOfflineKernelImpl(*MaybeKernelID, Self); + return nullptr; + } + std::shared_ptr tryGetKernel(detail::KernelNameStrRefT Name, const std::shared_ptr &Self) const { @@ -774,11 +785,7 @@ class kernel_bundle_impl { } // Fall back to regular offline compiled kernel_bundle look-up. - if (std::optional MaybeKernelID = - sycl::detail::ProgramManager::getInstance().tryGetSYCLKernelID( - Name)) - return tryGetOfflineKernel(*MaybeKernelID, Self); - return nullptr; + return tryGetOfflineKernel(Name, Self); } private: diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index b8b07e0053881..9a447c32a98b3 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2523,8 +2523,8 @@ getCGKernelInfo(const CGExecKernel &CommandGroup, ContextImplPtr ContextImpl, if (auto SyclKernelImpl = KernelBundleImplPtr - ? KernelBundleImplPtr->tryGetKernel(CommandGroup.MKernelName, - KernelBundleImplPtr) + ? KernelBundleImplPtr->tryGetOfflineKernel( + CommandGroup.MKernelName, KernelBundleImplPtr) : std::shared_ptr{nullptr}) { UrKernel = SyclKernelImpl->getHandleRef(); DeviceImageImpl = SyclKernelImpl->getDeviceImage(); diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 97e6e316fea72..60ee539e8fb70 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -445,10 +445,8 @@ event handler::finalize() { // Make sure implicit non-interop kernel bundles have the kernel if (!impl->isStateExplicitKernelBundle() && !(MKernel && MKernel->isInterop()) && - (KernelBundleImpPtr->empty() || - KernelBundleImpPtr->hasSYCLOfflineImages()) && - !KernelBundleImpPtr->tryGetKernel(MKernelName.data(), - KernelBundleImpPtr)) { + !KernelBundleImpPtr->tryGetOfflineKernel(MKernelName.data(), + KernelBundleImpPtr)) { auto Dev = impl->MGraph ? impl->MGraph->getDevice() : MQueue->get_device(); kernel_id KernelID =