[LIBOMPTARGET]Fix PR44933: fix crash because of the too early deinitialization of libomptarget.

[ABataev]

Summary

Instead of using global variables with unpredicted time of
deinitialization, use dynamically allocated variables with functions
explicitly marked as global constructor/destructor and priority. This
allows to prevent the crash because of the incorrect order of dynamic
libraries deinitialization.

[grokos]

I understand the rationale behind this patch, although I haven't been able to reproduce the problem. Maybe we are just being proactive here and stay on the safe side before any problem actually occurs. I'll keep the review open for a while just in case someone wants to raise any concerns, but to me it looks good.

[ABataev]

In D74837#1883416, @grokos wrote:

I understand the rationale behind this patch, although I haven't been able to reproduce the problem. Maybe we are just being proactive here and stay on the safe side before any problem actually occurs. I'll keep the review open for a while just in case someone wants to raise any concerns, but to me it looks good.

Did you try to run the test without this patch? For me, it crashes, just like the reproducer from the report

[grokos]

In D74837#1883449, @ABataev wrote:

Did you try to run the test without this patch? For me, it crashes, just like the reproducer from the report

Yes, I ran the provided test and for me it works correctly without the patch. But the nature of the problem makes reproducibility somewhat unpredictable, maybe on my system the RTL globals just happen to be deinitialized at the right time.

[ABataev]

Maybe it is just the problem of the test. I have CHECK-NOT in it to detect a possible crash. Maybe, just missed a variant for your system. Did you try the reproducer from the error report?

[grokos]

In D74837#1883479, @ABataev wrote:

Maybe it is just the problem of the test. I have CHECK-NOT in it to detect a possible crash. Maybe, just missed a variant for your system. Did you try the reproducer from the error report?

Yes, the reproducer works for me without this patch. But I'm targeting x86_64-pc-linux-gnu instead of nvptx64-nvidia-cuda, that might explain the difference in behavior (although I doubt it).

[ABataev]

In D74837#1883494, @grokos wrote:

In D74837#1883479, @ABataev wrote:

Maybe it is just the problem of the test. I have CHECK-NOT in it to detect a possible crash. Maybe, just missed a variant for your system. Did you try the reproducer from the error report?

Yes, the reproducer works for me without this patch. But I'm targeting x86_64-pc-linux-gnu instead of nvptx64-nvidia-cuda, that might explain the difference in behavior (although I doubt it).

No, I rather doubt in it. The test and the reproducer crash for me even when the host triple is specified as OpenMP target device.

[grokos]

I think we should move forward (I still haven't been able to reproduce the problem). LGTM.

openmp/libomptarget/src/api.cpp

 #include "private.h"
 #include "rtl.h"
 
 #include <climits>
 #include <cstring>
 #include <cstdlib>
 
 EXTERN int omp_get_num_devices(void) {
-  RTLsMtx.lock();
+  RTLsMtx->lock();
   size_t Devices_size = Devices.size();
-  RTLsMtx.unlock();
+  RTLsMtx->unlock();
 
   DP("Call to omp_get_num_devices returning %zd\n", Devices_size);
 
   return Devices_size;
 }
 
 EXTERN int omp_get_initial_device(void) {
   DP("Call to omp_get_initial_device returning %d\n", HOST_DEVICE);
if (!ptr) {
     return false;
   }
 
   if (device_num == omp_get_initial_device()) {
     DP("Call to omp_target_is_present on host, returning true\n");
     return true;
   }
 
-  RTLsMtx.lock();
+  RTLsMtx->lock();
   size_t Devices_size = Devices.size();
-  RTLsMtx.unlock();
+  RTLsMtx->unlock();
   if (Devices_size <= (size_t)device_num) {
     DP("Call to omp_target_is_present with invalid device ID, returning "
         "false\n");
     return false;
   }
 
   DeviceTy& Device = Devices[device_num];
   bool IsLast; // not used
   bool IsHostPtr;
   void *TgtPtr = Device.getTgtPtrBegin(ptr, 0, IsLast, false, IsHostPtr);
   int rc = (TgtPtr != NULL);
   // Under unified memory the host pointer can be returned by the
   // getTgtPtrBegin() function which means that there is no device
   // corresponding point for ptr. This function should return false
   // in that situation.
-  if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)
+  if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)
     rc = !IsHostPtr;
   DP("Call to omp_target_is_present returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_target_memcpy(void *dst, void *src, size_t length,
     size_t dst_offset, size_t src_offset, int dst_device, int src_device) {
   DP("Call to omp_target_memcpy, dst device %d, src device %d, "

openmp/libomptarget/src/device.cpp

if (lr.Flags.IsContained ||
     DP("Explicit extension of mapping is not allowed.\n");
   } else if (Size) {
     // If unified shared memory is active, implicitly mapped variables that are not
     // privatized use host address. Any explicitly mapped variables also use
     // host address where correctness is not impeded. In all other cases
     // maps are respected.
     // In addition to the mapping rules above, the close map
     // modifier forces the mapping of the variable to the device.
-    if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier) {
+    if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+        !HasCloseModifier) {
       DP("Return HstPtrBegin " DPxMOD " Size=%ld RefCount=%s\n",
          DPxPTR((uintptr_t)HstPtrBegin), Size, (UpdateRefCount ? " updated" : ""));
       IsHostPtr = true;
       rc = HstPtrBegin;
     } else {
       // If it is not contained and Size > 0 we should create a new entry for it.
       IsNew = true;
       uintptr_t tp = (uintptr_t)RTL->data_alloc(RTLDeviceID, Size, HstPtrBegin);
if (!IsLast && UpdateRefCount)
       HT.decRefCount();
 
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     DP("Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
         "Size=%ld,%s RefCount=%s\n", DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
         (UpdateRefCount ? " updated" : ""),
         HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
     rc = (void *)tp;
-  } else if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
+  } else if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
     // If the value isn't found in the mapping and unified shared memory
     // is on then it means we have stumbled upon a value which we need to
     // use directly from the host.
     DP("Get HstPtrBegin " DPxMOD " Size=%ld RefCount=%s\n",
        DPxPTR((uintptr_t)HstPtrBegin), Size, (UpdateRefCount ? " updated" : ""));
     IsHostPtr = true;
     rc = HstPtrBegin;
   }
if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     return (void *)tp;
   }
 
   return NULL;
 }
 
 int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
                             bool HasCloseModifier) {
-  if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier)
+  if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier)
     return OFFLOAD_SUCCESS;
   // Check if the pointer is contained in any sub-nodes.
   int rc;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     if (ForceDelete)
int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
   DataMapMtx.unlock();
   return rc;
 }
 
 /// Init device, should not be called directly.
 void DeviceTy::init() {
   // Make call to init_requires if it exists for this plugin.
   if (RTL->init_requires)
-    RTL->init_requires(RTLs.RequiresFlags);
+    RTL->init_requires(RTLs->RequiresFlags);
   int32_t rc = RTL->init_device(RTLDeviceID);
   if (rc == OFFLOAD_SUCCESS) {
     IsInit = true;
   }
 }
 
 /// Thread-safe method to initialize the device only once.
 int32_t DeviceTy::initOnce() {

 }
 
 /// Check whether a device has an associated RTL and initialize it if it's not
 /// already initialized.
 bool device_is_ready(int device_num) {
   DP("Checking whether device %d is ready.\n", device_num);
   // Devices.size() can only change while registering a new
   // library, so try to acquire the lock of RTLs' mutex.
-  RTLsMtx.lock();
+  RTLsMtx->lock();
   size_t Devices_size = Devices.size();
-  RTLsMtx.unlock();
+  RTLsMtx->unlock();
   if (Devices_size <= (size_t)device_num) {
     DP("Device ID  %d does not have a matching RTL\n", device_num);
     return false;
   }
 
   // Get device info
   DeviceTy &Device = Devices[device_num];
 

openmp/libomptarget/src/interface.cpp

case tgt_mandatory:
       }
       break;
   }
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags
 EXTERN void __tgt_register_requires(int64_t flags) {
-  RTLs.RegisterRequires(flags);
+  RTLs->RegisterRequires(flags);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds a target shared library to the target execution image
 EXTERN void __tgt_register_lib(__tgt_bin_desc *desc) {
-  RTLs.RegisterLib(desc);
+  RTLs->RegisterLib(desc);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// unloads a target shared library
 EXTERN void __tgt_unregister_lib(__tgt_bin_desc *desc) {
-  RTLs.UnregisterLib(desc);
+  RTLs->UnregisterLib(desc);
 }
 
 /// creates host-to-target data mapping, stores it in the
 /// libomptarget.so internal structure (an entry in a stack of data maps)
 /// and passes the data to the device.
 EXTERN void __tgt_target_data_begin(int64_t device_id, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types) {
   if (IsOffloadDisabled()) return;
EXTERN void __tgt_target_data_end(int64_t device_id, int32_t arg_num,
   if (IsOffloadDisabled()) return;
   DP("Entering data end region with %d mappings\n", arg_num);
 
   // No devices available?
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     device_id = omp_get_default_device();
   }
 
-  RTLsMtx.lock();
+  RTLsMtx->lock();
   size_t Devices_size = Devices.size();
-  RTLsMtx.unlock();
+  RTLsMtx->unlock();
   if (Devices_size <= (size_t)device_id) {
     DP("Device ID  %" PRId64 " does not have a matching RTL.\n", device_id);
     HandleTargetOutcome(false);
     return;
   }
 
   DeviceTy &Device = Devices[device_id];
   if (!Device.IsInit) {
EXTERN void __kmpc_push_target_tripcount(int64_t device_id,
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
     DP("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false);
     return;
   }
 
   DP("__kmpc_push_target_tripcount(%" PRId64 ", %" PRIu64 ")\n", device_id,
       loop_tripcount);
-  TblMapMtx.lock();
+  TblMapMtx->lock();
   Devices[device_id].LoopTripCnt.emplace(__kmpc_global_thread_num(NULL),
                                          loop_tripcount);
-  TblMapMtx.unlock();
+  TblMapMtx->unlock();
 }

openmp/libomptarget/src/omptarget.cpp

 static int InitLibrary(DeviceTy& Device) {
   /*
    * Map global data
    */
   int32_t device_id = Device.DeviceID;
   int rc = OFFLOAD_SUCCESS;
 
   Device.PendingGlobalsMtx.lock();
-  TrlTblMtx.lock();
+  TrlTblMtx->lock();
   for (HostEntriesBeginToTransTableTy::iterator
-      ii = HostEntriesBeginToTransTable.begin();
-      ii != HostEntriesBeginToTransTable.end(); ++ii) {
+      ii = HostEntriesBeginToTransTable->begin();
+      ii != HostEntriesBeginToTransTable->end(); ++ii) {
     TranslationTable *TransTable = &ii->second;
     if (TransTable->HostTable.EntriesBegin ==
         TransTable->HostTable.EntriesEnd) {
       // No host entry so no need to proceed
       continue;
     }
     if (TransTable->TargetsTable[device_id] != 0) {
       // Library entries have already been processed
for (__tgt_offload_entry *CurrDeviceEntry = TargetTable->EntriesBegin,
             (uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/,
             (uintptr_t)CurrHostEntry->addr + CurrHostEntry->size /*HstPtrEnd*/,
             (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/,
             true /*IsRefCountINF*/));
       }
     }
     Device.DataMapMtx.unlock();
   }
-  TrlTblMtx.unlock();
+  TrlTblMtx->unlock();
 
   if (rc != OFFLOAD_SUCCESS) {
     Device.PendingGlobalsMtx.unlock();
     return rc;
   }
 
   /*
    * Run ctors for static objects
if (arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
       uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
       void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
       DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
       args_base[i] = TgtPtrBase;
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
       bool copy = false;
-      if (!(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
+      if (!(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
           HasCloseModifier) {
         if (IsNew || (arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS)) {
           copy = true;
         } else if (arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) {
           // Copy data only if the "parent" struct has RefCount==1.
           int32_t parent_idx = member_of(arg_types[i]);
           uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
           assert(parent_rc > 0 && "parent struct not found");
if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
       DelEntry = false; // protect parent struct from being deallocated
     }
 
     if ((arg_types[i] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
       // Move data back to the host
       if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) {
         bool Always = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS;
         bool CopyMember = false;
-        if (!(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
+        if (!(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
             HasCloseModifier) {
           if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
               !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
             // Copy data only if the "parent" struct has RefCount==1.
             int32_t parent_idx = member_of(arg_types[i]);
             uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
             assert(parent_rc > 0 && "parent struct not found");
             if (parent_rc == 1) {
               CopyMember = true;
             }
           }
         }
 
         if ((DelEntry || Always || CopyMember) &&
-            !(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+            !(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
               TgtPtrBegin == HstPtrBegin)) {
           DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
               data_size, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
           int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, data_size);
           if (rt != OFFLOAD_SUCCESS) {
             DP("Copying data from device failed.\n");
             return OFFLOAD_FAIL;
           }
for (int32_t i = 0; i < arg_num; ++i) {
     bool IsLast, IsHostPtr;
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, MapSize, IsLast,
         false, IsHostPtr);
     if (!TgtPtrBegin) {
       DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
       continue;
     }
 
-    if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+    if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
         TgtPtrBegin == HstPtrBegin) {
       DP("hst data:" DPxMOD " unified and shared, becomes a noop\n",
          DPxPTR(HstPtrBegin));
       continue;
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) {
       DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",

 int target(int64_t device_id, void *host_ptr, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
     int32_t team_num, int32_t thread_limit, int IsTeamConstruct) {
   DeviceTy &Device = Devices[device_id];
 
   // Find the table information in the map or look it up in the translation
   // tables.
   TableMap *TM = 0;
-  TblMapMtx.lock();
-  HostPtrToTableMapTy::iterator TableMapIt = HostPtrToTableMap.find(host_ptr);
-  if (TableMapIt == HostPtrToTableMap.end()) {
+  TblMapMtx->lock();
+  HostPtrToTableMapTy::iterator TableMapIt = HostPtrToTableMap->find(host_ptr);
+  if (TableMapIt == HostPtrToTableMap->end()) {
     // We don't have a map. So search all the registered libraries.
-    TrlTblMtx.lock();
+    TrlTblMtx->lock();
     for (HostEntriesBeginToTransTableTy::iterator
-             ii = HostEntriesBeginToTransTable.begin(),
-             ie = HostEntriesBeginToTransTable.end();
+             ii = HostEntriesBeginToTransTable->begin(),
+             ie = HostEntriesBeginToTransTable->end();
          !TM && ii != ie; ++ii) {
       // get the translation table (which contains all the good info).
       TranslationTable *TransTable = &ii->second;
       // iterate over all the host table entries to see if we can locate the
       // host_ptr.
       __tgt_offload_entry *begin = TransTable->HostTable.EntriesBegin;
       __tgt_offload_entry *end = TransTable->HostTable.EntriesEnd;
       __tgt_offload_entry *cur = begin;
       for (uint32_t i = 0; cur < end; ++cur, ++i) {
         if (cur->addr != host_ptr)
           continue;
         // we got a match, now fill the HostPtrToTableMap so that we
         // may avoid this search next time.
-        TM = &HostPtrToTableMap[host_ptr];
+        TM = &(*HostPtrToTableMap)[host_ptr];
         TM->Table = TransTable;
         TM->Index = i;
         break;
       }
     }
-    TrlTblMtx.unlock();
+    TrlTblMtx->unlock();
   } else {
     TM = &TableMapIt->second;
   }
-  TblMapMtx.unlock();
+  TblMapMtx->unlock();
 
   // No map for this host pointer found!
   if (!TM) {
     DP("Host ptr " DPxMOD " does not have a matching target pointer.\n",
        DPxPTR(host_ptr));
     return OFFLOAD_FAIL;
   }
 
   // get target table.
-  TrlTblMtx.lock();
+  TrlTblMtx->lock();
   assert(TM->Table->TargetsTable.size() > (size_t)device_id &&
          "Not expecting a device ID outside the table's bounds!");
   __tgt_target_table *TargetTable = TM->Table->TargetsTable[device_id];
-  TrlTblMtx.unlock();
+  TrlTblMtx->unlock();
   assert(TargetTable && "Global data has not been mapped\n");
 
   // Move data to device.
   int rc = target_data_begin(Device, arg_num, args_base, args, arg_sizes,
       arg_types);
   if (rc != OFFLOAD_SUCCESS) {
     DP("Call to target_data_begin failed, abort target.\n");
     return OFFLOAD_FAIL;
if (!(arg_types[i] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
         void *Pointer_TgtPtrBegin =
             Device.getTgtPtrBegin(HstPtrVal, arg_sizes[i], IsLast, false,
                                   IsHostPtr);
         if (!Pointer_TgtPtrBegin) {
           DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
              DPxPTR(HstPtrVal));
           continue;
         }
-        if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+        if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
             TgtPtrBegin == HstPtrBegin) {
           DP("Unified memory is active, no need to map lambda captured"
              "variable (" DPxMOD ")\n", DPxPTR(HstPtrVal));
           continue;
         }
         DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n",
            DPxPTR(Pointer_TgtPtrBegin), DPxPTR(TgtPtrBegin));
         int rt = Device.data_submit(TgtPtrBegin, &Pointer_TgtPtrBegin,
#endif
     tgt_offsets.push_back(TgtBaseOffset);
   }
 
   assert(tgt_args.size() == tgt_offsets.size() &&
       "Size mismatch in arguments and offsets");
 
   // Pop loop trip count
   uint64_t ltc = 0;
-  TblMapMtx.lock();
+  TblMapMtx->lock();
   auto I = Device.LoopTripCnt.find(__kmpc_global_thread_num(NULL));
   if (I != Device.LoopTripCnt.end()) {
     ltc = I->second;
     Device.LoopTripCnt.erase(I);
     DP("loop trip count is %lu.\n", ltc);
   }
-  TblMapMtx.unlock();
+  TblMapMtx->unlock();
 
   // Launch device execution.
   DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
       TargetTable->EntriesBegin[TM->Index].name,
       DPxPTR(TargetTable->EntriesBegin[TM->Index].addr), TM->Index);
   if (IsTeamConstruct) {
     rc = Device.run_team_region(TargetTable->EntriesBegin[TM->Index].addr,
         &tgt_args[0], &tgt_offsets[0], tgt_args.size(), team_num,

openmp/libomptarget/src/rtl.h

public:
   void RegisterRequires(int64_t flags);
 
   // Register a shared library with all (compatible) RTLs.
   void RegisterLib(__tgt_bin_desc *desc);
 
   // Unregister a shared library from all RTLs.
   void UnregisterLib(__tgt_bin_desc *desc);
 };
-extern RTLsTy RTLs;
-extern std::mutex RTLsMtx;
+extern RTLsTy *RTLs;
+extern std::mutex *RTLsMtx;
 
 
 /// Map between the host entry begin and the translation table. Each
 /// registered library gets one TranslationTable. Use the map from
 /// __tgt_offload_entry so that we may quickly determine whether we
 /// are trying to (re)register an existing lib or really have a new one.
 struct TranslationTable {
   __tgt_target_table HostTable;
 
   // Image assigned to a given device.
   std::vector<__tgt_device_image *> TargetsImages; // One image per device ID.
 
   // Table of entry points or NULL if it was not already computed.
   std::vector<__tgt_target_table *> TargetsTable; // One table per device ID.
 };
 typedef std::map<__tgt_offload_entry *, TranslationTable>
     HostEntriesBeginToTransTableTy;
-extern HostEntriesBeginToTransTableTy HostEntriesBeginToTransTable;
-extern std::mutex TrlTblMtx;
+extern HostEntriesBeginToTransTableTy *HostEntriesBeginToTransTable;
+extern std::mutex *TrlTblMtx;
 
 /// Map between the host ptr and a table index
 struct TableMap {
   TranslationTable *Table; // table associated with the host ptr.
   uint32_t Index; // index in which the host ptr translated entry is found.
   TableMap() : Table(0), Index(0) {}
   TableMap(TranslationTable *table, uint32_t index)
       : Table(table), Index(index) {}
 };
 typedef std::map<void *, TableMap> HostPtrToTableMapTy;
-extern HostPtrToTableMapTy HostPtrToTableMap;
-extern std::mutex TblMapMtx;
+extern HostPtrToTableMapTy *HostPtrToTableMap;
+extern std::mutex *TblMapMtx;
 
 #endif

openmp/libomptarget/src/rtl.cpp

 
 // List of all plugins that can support offloading.
 static const char *RTLNames[] = {
     /* PowerPC target */ "libomptarget.rtl.ppc64.so",
     /* x86_64 target  */ "libomptarget.rtl.x86_64.so",
     /* CUDA target    */ "libomptarget.rtl.cuda.so",
     /* AArch64 target */ "libomptarget.rtl.aarch64.so"};
 
-RTLsTy RTLs;
-std::mutex RTLsMtx;
+RTLsTy *RTLs;
+std::mutex *RTLsMtx;
 
-HostEntriesBeginToTransTableTy HostEntriesBeginToTransTable;
-std::mutex TrlTblMtx;
+HostEntriesBeginToTransTableTy *HostEntriesBeginToTransTable;
+std::mutex *TrlTblMtx;
 
-HostPtrToTableMapTy HostPtrToTableMap;
-std::mutex TblMapMtx;
+HostPtrToTableMapTy *HostPtrToTableMap;
+std::mutex *TblMapMtx;
+
+__attribute__((constructor(0))) void init() {
+  DP("Init target library!\n");
+  RTLs = new RTLsTy();
+  RTLsMtx = new std::mutex();
+  HostEntriesBeginToTransTable = new HostEntriesBeginToTransTableTy();
+  TrlTblMtx = new std::mutex();
+  HostPtrToTableMap = new HostPtrToTableMapTy();
+  TblMapMtx = new std::mutex();
+}
+
+__attribute__((destructor(0))) void deinit() {
+  DP("Deinit target library!\n");
+  delete RTLs;
+  delete RTLsMtx;
+  delete HostEntriesBeginToTransTable;
+  delete TrlTblMtx;
+  delete HostPtrToTableMap;
+  delete TblMapMtx;
+}
 
 void RTLsTy::LoadRTLs() {
 #ifdef OMPTARGET_DEBUG
   if (char *envStr = getenv("LIBOMPTARGET_DEBUG")) {
     DebugLevel = std::stoi(envStr);
   }
 #endif // OMPTARGET_DEBUG
 
void RTLsTy::RegisterRequires(int64_t flags) {
   DP("New requires flags %ld compatible with existing %ld!\n",
      flags, RequiresFlags);
 }
 
 void RTLsTy::RegisterLib(__tgt_bin_desc *desc) {
   // Attempt to load all plugins available in the system.
   std::call_once(initFlag, &RTLsTy::LoadRTLs, this);
 
-  RTLsMtx.lock();
+  RTLsMtx->lock();
   // Register the images with the RTLs that understand them, if any.
   for (int32_t i = 0; i < desc->NumDeviceImages; ++i) {
     // Obtain the image.
     __tgt_device_image *img = &desc->DeviceImages[i];
 
     RTLInfoTy *FoundRTL = NULL;
 
     // Scan the RTLs that have associated images until we find one that supports
     // the current image.
-    for (auto &R : RTLs.AllRTLs) {
+    for (auto &R : AllRTLs) {
       if (!R.is_valid_binary(img)) {
         DP("Image " DPxMOD " is NOT compatible with RTL %s!\n",
             DPxPTR(img->ImageStart), R.RTLName.c_str());
         continue;
       }
 
       DP("Image " DPxMOD " is compatible with RTL %s!\n",
           DPxPTR(img->ImageStart), R.RTLName.c_str());
 
       // If this RTL is not already in use, initialize it.
       if (!R.isUsed) {
         // Initialize the device information for the RTL we are about to use.
         DeviceTy device(&R);
         size_t start = Devices.size();
         Devices.resize(start + R.NumberOfDevices, device);
         for (int32_t device_id = 0; device_id < R.NumberOfDevices;
             device_id++) {
           // global device ID
           Devices[start + device_id].DeviceID = start + device_id;
           // RTL local device ID
           Devices[start + device_id].RTLDeviceID = device_id;
         }
 
         // Initialize the index of this RTL and save it in the used RTLs.
-        R.Idx = (RTLs.UsedRTLs.empty())
+        R.Idx = (UsedRTLs.empty())
                     ? 0
-                    : RTLs.UsedRTLs.back()->Idx +
-                          RTLs.UsedRTLs.back()->NumberOfDevices;
+                    : UsedRTLs.back()->Idx + UsedRTLs.back()->NumberOfDevices;
         assert((size_t) R.Idx == start &&
             "RTL index should equal the number of devices used so far.");
         R.isUsed = true;
-        RTLs.UsedRTLs.push_back(&R);
+        UsedRTLs.push_back(&R);
 
         DP("RTL " DPxMOD " has index %d!\n", DPxPTR(R.LibraryHandler), R.Idx);
       }
 
       // Initialize (if necessary) translation table for this library.
-      TrlTblMtx.lock();
-      if(!HostEntriesBeginToTransTable.count(desc->HostEntriesBegin)){
+      TrlTblMtx->lock();
+      if(!HostEntriesBeginToTransTable->count(desc->HostEntriesBegin)){
         TranslationTable &tt =
-            HostEntriesBeginToTransTable[desc->HostEntriesBegin];
+            (*HostEntriesBeginToTransTable)[desc->HostEntriesBegin];
         tt.HostTable.EntriesBegin = desc->HostEntriesBegin;
         tt.HostTable.EntriesEnd = desc->HostEntriesEnd;
       }
 
       // Retrieve translation table for this library.
       TranslationTable &TransTable =
-          HostEntriesBeginToTransTable[desc->HostEntriesBegin];
+          (*HostEntriesBeginToTransTable)[desc->HostEntriesBegin];
 
       DP("Registering image " DPxMOD " with RTL %s!\n",
           DPxPTR(img->ImageStart), R.RTLName.c_str());
       RegisterImageIntoTranslationTable(TransTable, R, img);
-      TrlTblMtx.unlock();
+      TrlTblMtx->unlock();
       FoundRTL = &R;
 
       // Load ctors/dtors for static objects
       RegisterGlobalCtorsDtorsForImage(desc, img, FoundRTL);
 
       // if an RTL was found we are done - proceed to register the next image
       break;
     }
 
     if (!FoundRTL) {
       DP("No RTL found for image " DPxMOD "!\n", DPxPTR(img->ImageStart));
     }
   }
-  RTLsMtx.unlock();
+  RTLsMtx->unlock();
 
 
   DP("Done registering entries!\n");
 }
 
 void RTLsTy::UnregisterLib(__tgt_bin_desc *desc) {
   DP("Unloading target library!\n");
 
-  RTLsMtx.lock();
+  RTLsMtx->lock();
   // Find which RTL understands each image, if any.
   for (int32_t i = 0; i < desc->NumDeviceImages; ++i) {
     // Obtain the image.
     __tgt_device_image *img = &desc->DeviceImages[i];
 
     RTLInfoTy *FoundRTL = NULL;
 
     // Scan the RTLs that have associated images until we find one that supports
     // the current image. We only need to scan RTLs that are already being used.
-    for (auto *R : RTLs.UsedRTLs) {
+    for (auto *R : UsedRTLs) {
 
       assert(R->isUsed && "Expecting used RTLs.");
 
       if (!R->is_valid_binary(img)) {
         DP("Image " DPxMOD " is NOT compatible with RTL " DPxMOD "!\n",
             DPxPTR(img->ImageStart), DPxPTR(R->LibraryHandler));
         continue;
       }
for (int32_t i = 0; i < desc->NumDeviceImages; ++i) {
     }
 
     // if no RTL was found proceed to unregister the next image
     if (!FoundRTL){
       DP("No RTLs in use support the image " DPxMOD "!\n",
           DPxPTR(img->ImageStart));
     }
   }
-  RTLsMtx.unlock();
+  RTLsMtx->unlock();
   DP("Done unregistering images!\n");
 
   // Remove entries from HostPtrToTableMap
-  TblMapMtx.lock();
+  TblMapMtx->lock();
   for (__tgt_offload_entry *cur = desc->HostEntriesBegin;
       cur < desc->HostEntriesEnd; ++cur) {
-    HostPtrToTableMap.erase(cur->addr);
+    HostPtrToTableMap->erase(cur->addr);
   }
 
   // Remove translation table for this descriptor.
-  auto tt = HostEntriesBeginToTransTable.find(desc->HostEntriesBegin);
-  if (tt != HostEntriesBeginToTransTable.end()) {
+  auto tt = HostEntriesBeginToTransTable->find(desc->HostEntriesBegin);
+  if (tt != HostEntriesBeginToTransTable->end()) {
     DP("Removing translation table for descriptor " DPxMOD "\n",
         DPxPTR(desc->HostEntriesBegin));
-    HostEntriesBeginToTransTable.erase(tt);
+    HostEntriesBeginToTransTable->erase(tt);
   } else {
     DP("Translation table for descriptor " DPxMOD " cannot be found, probably "
         "it has been already removed.\n", DPxPTR(desc->HostEntriesBegin));
   }
 
-  TblMapMtx.unlock();
+  TblMapMtx->unlock();
 
   // TODO: Remove RTL and the devices it manages if it's not used anymore?
   // TODO: Write some RTL->unload_image(...) function?
 
   DP("Done unregistering library!\n");
 }

openmp/libomptarget/test/offloading/dynamic_module_load.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -DSHARED -shared -o %t.so && %clang %flags %s -o %t-aarch64-unknown-linux-gnu -ldl && %libomptarget-run-aarch64-unknown-linux-gnu %t.so 2>&1 | %fcheck-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -DSHARED -shared -o %t.so && %clang %flags %s -o %t-powerpc64-ibm-linux-gnu -ldl && %libomptarget-run-powerpc64-ibm-linux-gnu %t.so 2>&1 | %fcheck-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -DSHARED -shared -o %t.so && %clang %flags %s -o %t-powerpc64le-ibm-linux-gnu -ldl && %libomptarget-run-powerpc64le-ibm-linux-gnu %t.so 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -DSHARED -shared -o %t.so && %clang %flags %s -o %t-x86_64-pc-linux-gnu -ldl && %libomptarget-run-x86_64-pc-linux-gnu %t.so 2>&1 | %fcheck-x86_64-pc-linux-gnu
+
+#ifdef SHARED
+#include <stdio.h>
+int foo() {
+#pragma omp target
+  ;
+  printf("%s\n", "DONE.");
+  return 0;
+}
+#else
+#include <dlfcn.h>
+#include <stdio.h>
+int main(int argc, char **argv) {
+  void *Handle = dlopen(argv[1], RTLD_NOW);
+  int (*Foo)(void);
+
+  if (Handle == NULL) {
+    printf("dlopen() failed: %s\n", dlerror());
+    return 1;
+  }
+  Foo = (int (*)(void)) dlsym(Handle, "foo");
+  if (Handle == NULL) {
+    printf("dlsym() failed: %s\n", dlerror());
+    return 1;
+  }
+  // CHECK: DONE.
+  // CHECK-NOT: {{abort|fault}}
+  return Foo();
+}
+#endif