diff --git a/openmp/libomptarget/include/PluginManager.h b/openmp/libomptarget/include/PluginManager.h index ec5d98dc8cd30..5e5306ac776f0 100644 --- a/openmp/libomptarget/include/PluginManager.h +++ b/openmp/libomptarget/include/PluginManager.h @@ -206,6 +206,12 @@ struct PluginManager { ProtectedObj Devices; }; +/// Initialize the plugin manager and OpenMP runtime. +void initRuntime(); + +/// Deinitialize the plugin and delete it. +void deinitRuntime(); + extern PluginManager *PM; #endif // OMPTARGET_PLUGIN_MANAGER_H diff --git a/openmp/libomptarget/include/omptarget.h b/openmp/libomptarget/include/omptarget.h index c4faa23427f11..9a2bd1340e3b4 100644 --- a/openmp/libomptarget/include/omptarget.h +++ b/openmp/libomptarget/include/omptarget.h @@ -312,6 +312,12 @@ void *llvm_omp_target_dynamic_shared_alloc(); /// add the clauses of the requires directives in a given file void __tgt_register_requires(int64_t Flags); +/// Initializes the runtime library. +void __tgt_rtl_init(); + +/// Deinitializes the runtime library. +void __tgt_rtl_deinit(); + /// adds a target shared library to the target execution image void __tgt_register_lib(__tgt_bin_desc *Desc); diff --git a/openmp/libomptarget/src/OffloadRTL.cpp b/openmp/libomptarget/src/OffloadRTL.cpp index 86ef0d5bc91cf..dd75b1b181505 100644 --- a/openmp/libomptarget/src/OffloadRTL.cpp +++ b/openmp/libomptarget/src/OffloadRTL.cpp @@ -20,25 +20,39 @@ extern void llvm::omp::target::ompt::connectLibrary(); #endif -__attribute__((constructor(101))) void init() { +static std::mutex PluginMtx; +static uint32_t RefCount = 0; + +void initRuntime() { + std::scoped_lock Lock(PluginMtx); Profiler::get(); TIMESCOPE(); - DP("Init offload library!\n"); - - PM = new PluginManager(); + if (PM == nullptr) + PM = new PluginManager(); + RefCount++; + if (RefCount == 1) { + DP("Init offload library!\n"); #ifdef OMPT_SUPPORT - // Initialize OMPT first - llvm::omp::target::ompt::connectLibrary(); + // Initialize OMPT first + llvm::omp::target::ompt::connectLibrary(); #endif - PM->init(); - - PM->registerDelayedLibraries(); + PM->init(); + PM->registerDelayedLibraries(); + } } -__attribute__((destructor(101))) void deinit() { - DP("Deinit offload library!\n"); - delete PM; +void deinitRuntime() { + std::scoped_lock Lock(PluginMtx); + assert(PM && "Runtime not initialized"); + + if (RefCount == 1) { + DP("Deinit offload library!\n"); + delete PM; + PM = nullptr; + } + + RefCount--; } diff --git a/openmp/libomptarget/src/PluginManager.cpp b/openmp/libomptarget/src/PluginManager.cpp index 34f1f4969da30..09f9c6400569c 100644 --- a/openmp/libomptarget/src/PluginManager.cpp +++ b/openmp/libomptarget/src/PluginManager.cpp @@ -21,7 +21,7 @@ using namespace llvm; using namespace llvm::sys; -PluginManager *PM; +PluginManager *PM = nullptr; // List of all plugins that can support offloading. static const char *RTLNames[] = {ENABLED_OFFLOAD_PLUGINS}; diff --git a/openmp/libomptarget/src/exports b/openmp/libomptarget/src/exports index af882a2642647..d5432a9eed380 100644 --- a/openmp/libomptarget/src/exports +++ b/openmp/libomptarget/src/exports @@ -1,5 +1,7 @@ VERS1.0 { global: + __tgt_rtl_init; + __tgt_rtl_deinit; __tgt_register_requires; __tgt_register_lib; __tgt_unregister_lib; diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp index d2707f39a1aa3..8b89bc3ff7124 100644 --- a/openmp/libomptarget/src/interface.cpp +++ b/openmp/libomptarget/src/interface.cpp @@ -38,9 +38,13 @@ EXTERN void __tgt_register_requires(int64_t Flags) { __PRETTY_FUNCTION__); } +EXTERN void __tgt_rtl_init() { initRuntime(); } +EXTERN void __tgt_rtl_deinit() { deinitRuntime(); } + //////////////////////////////////////////////////////////////////////////////// /// adds a target shared library to the target execution image EXTERN void __tgt_register_lib(__tgt_bin_desc *Desc) { + initRuntime(); if (PM->delayRegisterLib(Desc)) return; @@ -49,12 +53,17 @@ EXTERN void __tgt_register_lib(__tgt_bin_desc *Desc) { //////////////////////////////////////////////////////////////////////////////// /// Initialize all available devices without registering any image -EXTERN void __tgt_init_all_rtls() { PM->initAllPlugins(); } +EXTERN void __tgt_init_all_rtls() { + assert(PM && "Runtime not initialized"); + PM->initAllPlugins(); +} //////////////////////////////////////////////////////////////////////////////// /// unloads a target shared library EXTERN void __tgt_unregister_lib(__tgt_bin_desc *Desc) { PM->unregisterLib(Desc); + + deinitRuntime(); } template @@ -64,6 +73,7 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase, map_var_info_t *ArgNames, void **ArgMappers, TargetDataFuncPtrTy TargetDataFunction, const char *RegionTypeMsg, const char *RegionName) { + assert(PM && "Runtime not initialized"); static_assert(std::is_convertible_v, "TargetAsyncInfoTy must be convertible to AsyncInfoTy."); @@ -239,6 +249,7 @@ template static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams, int32_t ThreadLimit, void *HostPtr, KernelArgsTy *KernelArgs) { + assert(PM && "Runtime not initialized"); static_assert(std::is_convertible_v, "Target AsyncInfoTy must be convertible to AsyncInfoTy."); DP("Entering target region for device %" PRId64 " with entry point " DPxMOD @@ -345,6 +356,7 @@ EXTERN int __tgt_activate_record_replay(int64_t DeviceId, uint64_t MemorySize, void *VAddr, bool IsRecord, bool SaveOutput, uint64_t &ReqPtrArgOffset) { + assert(PM && "Runtime not initialized"); OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0))); auto DeviceOrErr = PM->getDevice(DeviceId); if (!DeviceOrErr) @@ -380,7 +392,7 @@ EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId, ptrdiff_t *TgtOffsets, int32_t NumArgs, int32_t NumTeams, int32_t ThreadLimit, uint64_t LoopTripCount) { - + assert(PM && "Runtime not initialized"); OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0))); if (checkDeviceAndCtors(DeviceId, Loc)) { DP("Not offloading to device %" PRId64 "\n", DeviceId); @@ -431,6 +443,7 @@ EXTERN void __tgt_push_mapper_component(void *RtMapperHandle, void *Base, } EXTERN void __tgt_set_info_flag(uint32_t NewInfoLevel) { + assert(PM && "Runtime not initialized"); std::atomic &InfoLevel = getInfoLevelInternal(); InfoLevel.store(NewInfoLevel); for (auto &R : PM->pluginAdaptors()) { @@ -440,6 +453,7 @@ EXTERN void __tgt_set_info_flag(uint32_t NewInfoLevel) { } EXTERN int __tgt_print_device_info(int64_t DeviceId) { + assert(PM && "Runtime not initialized"); auto DeviceOrErr = PM->getDevice(DeviceId); if (!DeviceOrErr) FATAL_MESSAGE(DeviceId, "%s", toString(DeviceOrErr.takeError()).c_str()); @@ -448,7 +462,9 @@ EXTERN int __tgt_print_device_info(int64_t DeviceId) { } EXTERN void __tgt_target_nowait_query(void **AsyncHandle) { + assert(PM && "Runtime not initialized"); OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0))); + if (!AsyncHandle || !*AsyncHandle) { FATAL_MESSAGE0( 1, "Receive an invalid async handle from the current OpenMP task. Is " diff --git a/openmp/libomptarget/test/offloading/runtime_init.c b/openmp/libomptarget/test/offloading/runtime_init.c new file mode 100644 index 0000000000000..96fd50f51da1e --- /dev/null +++ b/openmp/libomptarget/test/offloading/runtime_init.c @@ -0,0 +1,30 @@ +// RUN: %libomptarget-compile-generic +// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \ +// RUN: %fcheck-generic + +// REQUIRES: libomptarget-debug + +#include +#include + +extern void __tgt_rtl_init(void); +extern void __tgt_rtl_deinit(void); + +// Sanity checks to make sure that this works and is thread safe. +int main() { + // CHECK: Init offload library! + // CHECK: Deinit offload library! + __tgt_rtl_init(); +#pragma omp parallel num_threads(8) + { + __tgt_rtl_init(); + __tgt_rtl_deinit(); + } + __tgt_rtl_deinit(); + + __tgt_rtl_init(); + __tgt_rtl_deinit(); + + // CHECK: PASS + printf("PASS\n"); +}