From ac9064b729f92bb2e779b5896137cda7527a1c08 Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Fri, 23 Oct 2026 11:36:03 -0700 Subject: [PATCH 1/2] [SYCL] Keep kernel compiler resources alive during static destruction Kernel bundles built from SYCL source can be stored in global or static caches and released after main() returns. Before: - Destroying a cached source-built bundle at process exit could abort in ProgramManager::removeImages() because its image tables had already been freed by the runtime jit_compiler destructor. - The runtime jit_compiler and the JIT library's JITContext were function-local static objects. When initialized after the user's cache, they could be destroyed first, leaving bundle cleanup with dangling image descriptors and an already-destroyed binary registry. After: - GlobalHandler owns the runtime jit_compiler and releases it after ProgramManager during late runtime shutdown, keeping image tables valid throughout user static destruction. - On non-Windows platforms, allocate JITContext lazily on the heap and delete it from a late JIT-library unload handler instead of ordinary static destruction. Cached bundles can unregister their images and destroy their code blobs while both owners are still alive. Add a regression test for both ordinary bundle destruction and cleanup from a static cache. Executing cached kernels in the cache destructor provides an additional check that the resources remain usable. Assisted-by: OpenCode (OpenAI GPT-6.1) --- .../lib/translation/JITContext.cpp | 14 ++++ .../jit-compiler/lib/translation/JITContext.h | 14 ++++ sycl/source/detail/global_handler.cpp | 13 ++++ sycl/source/detail/global_handler.hpp | 3 + sycl/source/detail/jit_compiler.cpp | 5 ++ sycl/source/detail/jit_compiler.hpp | 8 +-- .../KernelCompiler/sycl_static_lifetime.cpp | 64 +++++++++++++++++++ 7 files changed, 117 insertions(+), 4 deletions(-) create mode 100644 sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp diff --git a/sycl-jit/jit-compiler/lib/translation/JITContext.cpp b/sycl-jit/jit-compiler/lib/translation/JITContext.cpp index a06c6ff81c65b..d7af4cc32ba0b 100644 --- a/sycl-jit/jit-compiler/lib/translation/JITContext.cpp +++ b/sycl-jit/jit-compiler/lib/translation/JITContext.cpp @@ -10,6 +10,20 @@ using namespace jit_compiler; +#ifndef _WIN32 +JITContext *JITContext::Instance = nullptr; + +void JITContext::destroyInstance() { + delete Instance; + Instance = nullptr; +} + +// Match the runtime's late unload priority, after user static destructors. +__attribute__((destructor(110))) static void jitUnload() { + JITContext::destroyInstance(); +} +#endif + JITBinary::JITBinary(std::string &&Binary, BinaryFormat Fmt) : Blob{std::move(Binary)}, Format{Fmt} {} diff --git a/sycl-jit/jit-compiler/lib/translation/JITContext.h b/sycl-jit/jit-compiler/lib/translation/JITContext.h index 83e08efd06109..9acdebc429bf4 100644 --- a/sycl-jit/jit-compiler/lib/translation/JITContext.h +++ b/sycl-jit/jit-compiler/lib/translation/JITContext.h @@ -51,10 +51,21 @@ class JITContext { public: static JITContext &getInstance() { +#ifdef _WIN32 static JITContext Instance{}; return Instance; +#else + // Runtime-compiled bundles can outlive ordinary static destructors. Keep + // their binary storage alive until the JIT library's late unload handler. + static JITContext &Context = *(Instance = new JITContext{}); + return Context; +#endif } +#ifndef _WIN32 + static void destroyInstance(); +#endif + template JITBinary &emplaceBinary(Ts &&...Args) { WriteLockT WriteLock{BinariesMutex}; auto JBUPtr = std::make_unique(std::forward(Args)...); @@ -69,6 +80,9 @@ class JITContext { } private: +#ifndef _WIN32 + static JITContext *Instance; +#endif JITContext() = default; ~JITContext() = default; JITContext(const JITContext &) = delete; diff --git a/sycl/source/detail/global_handler.cpp b/sycl/source/detail/global_handler.cpp index 0b7ac8bc9e8db..4ee4123a1bcff 100644 --- a/sycl/source/detail/global_handler.cpp +++ b/sycl/source/detail/global_handler.cpp @@ -15,6 +15,7 @@ #include #include #include +#include #include #include #include @@ -206,6 +207,15 @@ ProgramManager &GlobalHandler::getProgramManager() { return PM; } +#if SYCL_EXT_JIT_ENABLE +jit_compiler &GlobalHandler::getJITCompiler() { + const LockGuard Lock{MJITCompiler.Lock}; + if (!MJITCompiler.Inst) + MJITCompiler.Inst.reset(new jit_compiler()); + return *MJITCompiler.Inst; +} +#endif + std::unordered_map> & GlobalHandler::getPlatformToDefaultContextCache() { // The optimization with static reference is not done because @@ -447,6 +457,9 @@ void shutdown_late() { GlobalHandler::RTGlobalObjHandler->MScheduler.Inst.reset(nullptr); GlobalHandler::RTGlobalObjHandler->MProgramManager.Inst.reset(nullptr); + // The JIT compiler owns image tables referenced by the program manager. + GlobalHandler::RTGlobalObjHandler->MJITCompiler.Inst.reset(nullptr); + // Clear the adapters and reset the instance if it was there. GlobalHandler::RTGlobalObjHandler->unloadAdapters(); if (GlobalHandler::RTGlobalObjHandler->MAdapters.Inst) diff --git a/sycl/source/detail/global_handler.hpp b/sycl/source/detail/global_handler.hpp index 9c639a00cf2b4..26566bb330d7d 100644 --- a/sycl/source/detail/global_handler.hpp +++ b/sycl/source/detail/global_handler.hpp @@ -23,6 +23,7 @@ class platform_impl; class context_impl; class Scheduler; class ProgramManager; +class jit_compiler; class Sync; class adapter_impl; class ods_target_list; @@ -77,6 +78,7 @@ class GlobalHandler { /// returned access object keeps the Scheduler alive until it is destroyed. static SchedulerAccess getSchedulerAccess() noexcept; ProgramManager &getProgramManager(); + jit_compiler &getJITCompiler(); Sync &getSync(); std::vector> &getPlatformCache(); @@ -168,6 +170,7 @@ class GlobalHandler { // count accesses that must finish before the Scheduler can be destroyed. static std::atomic MSchedulerAccessState; InstWithLock MProgramManager; + InstWithLock MJITCompiler; InstWithLock MSync; InstWithLock>> MPlatformCache; InstWithLock< diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index c1227e3b4d141..4409c5f66f15d 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include #if SYCL_EXT_JIT_ENABLE +#include #include #include #include @@ -16,6 +17,10 @@ namespace sycl { inline namespace _V1 { namespace detail { +jit_compiler &jit_compiler::get_instance() { + return GlobalHandler::instance().getJITCompiler(); +} + std::function jit_compiler::CustomDeleterForLibHandle = [](void *StoredPtr) { if (!StoredPtr) diff --git a/sycl/source/detail/jit_compiler.hpp b/sycl/source/detail/jit_compiler.hpp index e37c9e4beb20a..d415b63a384cb 100644 --- a/sycl/source/detail/jit_compiler.hpp +++ b/sycl/source/detail/jit_compiler.hpp @@ -54,10 +54,7 @@ class jit_compiler { bool isAvailable() { return Available; } - static jit_compiler &get_instance() { - static jit_compiler instance{}; - return instance; - } + static jit_compiler &get_instance(); private: jit_compiler(); @@ -67,6 +64,9 @@ class jit_compiler { jit_compiler &operator=(const jit_compiler &) = delete; jit_compiler &operator=(const jit_compiler &&) = delete; + friend class GlobalHandler; + friend struct std::default_delete; + sycl_device_binaries createDeviceBinaries(const ::jit_compiler::RTCBundleInfo &BundleInfo, const std::string &Prefix); diff --git a/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp b/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp new file mode 100644 index 0000000000000..a2c27f8bbabdf --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp @@ -0,0 +1,64 @@ +// REQUIRES: sycl-jit, aspect-usm_shared_allocations + +// RUN: %{build} -o %t.out +// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out +// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out keep + +// Cached source-built bundles must support execution and cleanup during static +// destruction, after ordinary runtime static destructors would have run. +#include +#include +#include + +#include +#include +#include + +namespace syclex = sycl::ext::oneapi::experimental; +using ExecutableBundle = sycl::kernel_bundle; + +constexpr const char *Source = R"( +#include +#include +extern "C" SYCL_EXTERNAL +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::single_task_kernel)) +void increment(int *Value) { ++*Value; } +)"; + +struct KernelCache { + std::vector Bundles; + std::optional Queue; + + ~KernelCache() { + for (auto &Bundle : Bundles) { + auto Kernel = Bundle.ext_oneapi_get_kernel("increment"); + assert(Kernel.get_info() == 1); + int *Value = sycl::malloc_shared(1, *Queue); + *Value = 0; + Queue->submit([&](sycl::handler &CGH) { + CGH.set_arg(0, Value); + CGH.single_task(Kernel); + }); + Queue->wait_and_throw(); + assert(*Value == 1); + sycl::free(Value, *Queue); + } + } +}; + +int main(int Argc, char **) { + sycl::queue Q; + static KernelCache Cache; + assert(Q.get_device().ext_oneapi_can_build(syclex::source_language::sycl)); + auto SourceBundle = syclex::create_kernel_bundle_from_source( + Q.get_context(), syclex::source_language::sycl, Source); + for (int I = 0; I < 2; ++I) { + auto Bundle = syclex::build(SourceBundle); + assert(Bundle.ext_oneapi_get_kernel("increment") + .get_info() == 1); + if (Argc > 1) { + Cache.Queue = Q; + Cache.Bundles.push_back(Bundle); + } + } +} From de3d182444efc27e35210084b118c5ef94b06efe Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Tue, 27 Oct 2026 09:32:55 -0700 Subject: [PATCH 2/2] [SYCL][E2E] Fix test --- sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp b/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp index a2c27f8bbabdf..9555f3b652e54 100644 --- a/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp +++ b/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp @@ -1,4 +1,6 @@ // REQUIRES: sycl-jit, aspect-usm_shared_allocations +// UNSUPPORTED: target-native_cpu +// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20142 // RUN: %{build} -o %t.out // RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out @@ -32,7 +34,6 @@ struct KernelCache { ~KernelCache() { for (auto &Bundle : Bundles) { auto Kernel = Bundle.ext_oneapi_get_kernel("increment"); - assert(Kernel.get_info() == 1); int *Value = sycl::malloc_shared(1, *Queue); *Value = 0; Queue->submit([&](sycl::handler &CGH) { @@ -54,8 +55,8 @@ int main(int Argc, char **) { Q.get_context(), syclex::source_language::sycl, Source); for (int I = 0; I < 2; ++I) { auto Bundle = syclex::build(SourceBundle); - assert(Bundle.ext_oneapi_get_kernel("increment") - .get_info() == 1); + assert(Bundle.ext_oneapi_has_kernel("increment")); + auto Kernel = Bundle.ext_oneapi_get_kernel("increment"); if (Argc > 1) { Cache.Queue = Q; Cache.Bundles.push_back(Bundle);