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..9555f3b652e54 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/sycl_static_lifetime.cpp @@ -0,0 +1,65 @@ +// 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 +// 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"); + 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_has_kernel("increment")); + auto Kernel = Bundle.ext_oneapi_get_kernel("increment"); + if (Argc > 1) { + Cache.Queue = Q; + Cache.Bundles.push_back(Bundle); + } + } +}