[Libomptarget] Support device destructors more correctly

[jhuber6]

Summary

Previous patched addressed problems with the plugin's teardown order by
implementing a specific function used to initialize and uninitialize the
plugins. This allows us to more correctly control the teardown order.
This is important for handling device destructors which were previously
problematic as the runtime may have been unititialized prior to
attempting to run the destructors. This patch adds support for the new
init / deinit plugin functions to the CUDA and generic plugins. We can
now correctly test device destructors.

[jhuber6]

This seems to cause problems for CUDA, we get a lot of errors about being unable to destroy a stream or an event.

openmp/libomptarget/plugins/cuda/src/rtl.cpp

int setContext(int DeviceId) {
     CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
     if (!checkResult(Err, "error returned from cuCtxSetCurrent"))
       return OFFLOAD_FAIL;
 
     return OFFLOAD_SUCCESS;
   }
 };
 
-DeviceRTLTy DeviceRTL;
+DeviceRTLTy *DeviceRTL = nullptr;
 } // namespace
 
 // Exposed library API function
 #ifdef __cplusplus
 extern "C" {
 #endif
 
+int32_t __tgt_rtl_init_plugin() {
+  DeviceRTL = new DeviceRTLTy;
+  return DeviceRTL ? OFFLOAD_SUCCESS : OFFLOAD_FAIL;
+}
+
+int32_t __tgt_rtl_deinit_plugin() {
+  if (DeviceRTL)
+    delete DeviceRTL;
+  return OFFLOAD_SUCCESS;
+}
+
 int32_t __tgt_rtl_is_valid_binary(__tgt_device_image *Image) {
   return elf_check_machine(Image, /* EM_CUDA */ 190);
 }
 
 int32_t __tgt_rtl_is_valid_binary_info(__tgt_device_image *image,
                                        __tgt_image_info *info) {
   if (!__tgt_rtl_is_valid_binary(image))
     return false;
for (int32_t DeviceId = 0; DeviceId < NumberOfDevices; ++DeviceId) {
     if (ArchStr != info->Arch)
       return false;
   }
 
   DP("Image has compatible compute capability: %s\n", info->Arch);
   return true;
 }
 
-int32_t __tgt_rtl_number_of_devices() { return DeviceRTL.getNumOfDevices(); }
+int32_t __tgt_rtl_number_of_devices() { return DeviceRTL->getNumOfDevices(); }
 
 int64_t __tgt_rtl_init_requires(int64_t RequiresFlags) {
   DP("Init requires flags to %" PRId64 "\n", RequiresFlags);
-  DeviceRTL.setRequiresFlag(RequiresFlags);
+  DeviceRTL->setRequiresFlag(RequiresFlags);
   return RequiresFlags;
 }
 
 int32_t __tgt_rtl_is_data_exchangable(int32_t SrcDevId, int DstDevId) {
-  if (DeviceRTL.isValidDeviceId(SrcDevId) &&
-      DeviceRTL.isValidDeviceId(DstDevId))
+  if (DeviceRTL->isValidDeviceId(SrcDevId) &&
+      DeviceRTL->isValidDeviceId(DstDevId))
     return 1;
 
   return 0;
 }
 
 int32_t __tgt_rtl_init_device(int32_t DeviceId) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // Context is set when init the device.
 
-  return DeviceRTL.initDevice(DeviceId);
+  return DeviceRTL->initDevice(DeviceId);
 }
 
 int32_t __tgt_rtl_deinit_device(int32_t DeviceId) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // Context is set when deinit the device.
 
-  return DeviceRTL.deinitDevice(DeviceId);
+  return DeviceRTL->deinitDevice(DeviceId);
 }
 
 __tgt_target_table *__tgt_rtl_load_binary(int32_t DeviceId,
                                           __tgt_device_image *Image) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return nullptr;
 
-  return DeviceRTL.loadBinary(DeviceId, Image);
+  return DeviceRTL->loadBinary(DeviceId, Image);
 }
 
 void *__tgt_rtl_data_alloc(int32_t DeviceId, int64_t Size, void *,
                            int32_t Kind) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return nullptr;
 
-  return DeviceRTL.dataAlloc(DeviceId, Size, (TargetAllocTy)Kind);
+  return DeviceRTL->dataAlloc(DeviceId, Size, (TargetAllocTy)Kind);
 }
 
 int32_t __tgt_rtl_data_submit(int32_t DeviceId, void *TgtPtr, void *HstPtr,
                               int64_t Size) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // Context is set in __tgt_rtl_data_submit_async.
 
   __tgt_async_info AsyncInfo;
   const int32_t Rc =
       __tgt_rtl_data_submit_async(DeviceId, TgtPtr, HstPtr, Size, &AsyncInfo);
   if (Rc != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return __tgt_rtl_synchronize(DeviceId, &AsyncInfo);
 }
 
 int32_t __tgt_rtl_data_submit_async(int32_t DeviceId, void *TgtPtr,
                                     void *HstPtr, int64_t Size,
                                     __tgt_async_info *AsyncInfoPtr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   assert(AsyncInfoPtr && "async_info_ptr is nullptr");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.dataSubmit(DeviceId, TgtPtr, HstPtr, Size, AsyncInfoPtr);
+  return DeviceRTL->dataSubmit(DeviceId, TgtPtr, HstPtr, Size, AsyncInfoPtr);
 }
 
 int32_t __tgt_rtl_data_retrieve(int32_t DeviceId, void *HstPtr, void *TgtPtr,
                                 int64_t Size) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // Context is set in __tgt_rtl_data_retrieve_async.
 
   __tgt_async_info AsyncInfo;
   const int32_t Rc =
       __tgt_rtl_data_retrieve_async(DeviceId, HstPtr, TgtPtr, Size, &AsyncInfo);
   if (Rc != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return __tgt_rtl_synchronize(DeviceId, &AsyncInfo);
 }
 
 int32_t __tgt_rtl_data_retrieve_async(int32_t DeviceId, void *HstPtr,
                                       void *TgtPtr, int64_t Size,
                                       __tgt_async_info *AsyncInfoPtr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   assert(AsyncInfoPtr && "async_info_ptr is nullptr");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.dataRetrieve(DeviceId, HstPtr, TgtPtr, Size, AsyncInfoPtr);
+  return DeviceRTL->dataRetrieve(DeviceId, HstPtr, TgtPtr, Size, AsyncInfoPtr);
 }
 
 int32_t __tgt_rtl_data_exchange_async(int32_t SrcDevId, void *SrcPtr,
                                       int DstDevId, void *DstPtr, int64_t Size,
                                       __tgt_async_info *AsyncInfo) {
-  assert(DeviceRTL.isValidDeviceId(SrcDevId) && "src_dev_id is invalid");
-  assert(DeviceRTL.isValidDeviceId(DstDevId) && "dst_dev_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(SrcDevId) && "src_dev_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DstDevId) && "dst_dev_id is invalid");
   assert(AsyncInfo && "AsyncInfo is nullptr");
 
-  if (DeviceRTL.setContext(SrcDevId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(SrcDevId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.dataExchange(SrcDevId, SrcPtr, DstDevId, DstPtr, Size,
+  return DeviceRTL->dataExchange(SrcDevId, SrcPtr, DstDevId, DstPtr, Size,
                                  AsyncInfo);
 }
 
 int32_t __tgt_rtl_data_exchange(int32_t SrcDevId, void *SrcPtr,
                                 int32_t DstDevId, void *DstPtr, int64_t Size) {
-  assert(DeviceRTL.isValidDeviceId(SrcDevId) && "src_dev_id is invalid");
-  assert(DeviceRTL.isValidDeviceId(DstDevId) && "dst_dev_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(SrcDevId) && "src_dev_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DstDevId) && "dst_dev_id is invalid");
   // Context is set in __tgt_rtl_data_exchange_async.
 
   __tgt_async_info AsyncInfo;
   const int32_t Rc = __tgt_rtl_data_exchange_async(SrcDevId, SrcPtr, DstDevId,
                                                    DstPtr, Size, &AsyncInfo);
   if (Rc != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return __tgt_rtl_synchronize(SrcDevId, &AsyncInfo);
 }
 
 int32_t __tgt_rtl_data_delete(int32_t DeviceId, void *TgtPtr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.dataDelete(DeviceId, TgtPtr);
+  return DeviceRTL->dataDelete(DeviceId, TgtPtr);
 }
 
 int32_t __tgt_rtl_run_target_team_region(int32_t DeviceId, void *TgtEntryPtr,
                                          void **TgtArgs, ptrdiff_t *TgtOffsets,
                                          int32_t ArgNum, int32_t TeamNum,
                                          int32_t ThreadLimit,
                                          uint64_t LoopTripcount) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // Context is set in __tgt_rtl_run_target_team_region_async.
 
   __tgt_async_info AsyncInfo;
   const int32_t Rc = __tgt_rtl_run_target_team_region_async(
       DeviceId, TgtEntryPtr, TgtArgs, TgtOffsets, ArgNum, TeamNum, ThreadLimit,
       LoopTripcount, &AsyncInfo);
   if (Rc != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return __tgt_rtl_synchronize(DeviceId, &AsyncInfo);
 }
 
 int32_t __tgt_rtl_run_target_team_region_async(
     int32_t DeviceId, void *TgtEntryPtr, void **TgtArgs, ptrdiff_t *TgtOffsets,
     int32_t ArgNum, int32_t TeamNum, int32_t ThreadLimit,
     uint64_t LoopTripcount, __tgt_async_info *AsyncInfoPtr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.runTargetTeamRegion(DeviceId, TgtEntryPtr, TgtArgs,
-                                       TgtOffsets, ArgNum, TeamNum, ThreadLimit,
-                                       LoopTripcount, AsyncInfoPtr);
+  return DeviceRTL->runTargetTeamRegion(
+      DeviceId, TgtEntryPtr, TgtArgs, TgtOffsets, ArgNum, TeamNum, ThreadLimit,
 }
 
 int32_t __tgt_rtl_run_target_region(int32_t DeviceId, void *TgtEntryPtr,
                                     void **TgtArgs, ptrdiff_t *TgtOffsets,
                                     int32_t ArgNum) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // Context is set in __tgt_rtl_run_target_region_async.
 
   __tgt_async_info AsyncInfo;
   const int32_t Rc = __tgt_rtl_run_target_region_async(
       DeviceId, TgtEntryPtr, TgtArgs, TgtOffsets, ArgNum, &AsyncInfo);
   if (Rc != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return __tgt_rtl_synchronize(DeviceId, &AsyncInfo);
 }
 
 int32_t __tgt_rtl_run_target_region_async(int32_t DeviceId, void *TgtEntryPtr,
                                           void **TgtArgs, ptrdiff_t *TgtOffsets,
                                           int32_t ArgNum,
                                           __tgt_async_info *AsyncInfoPtr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // Context is set in __tgt_rtl_run_target_team_region_async.
   return __tgt_rtl_run_target_team_region_async(
       DeviceId, TgtEntryPtr, TgtArgs, TgtOffsets, ArgNum,
       /* team num*/ 1, /* thread_limit */ 1, /* loop_tripcount */ 0,
       AsyncInfoPtr);
 }
 
 int32_t __tgt_rtl_synchronize(int32_t DeviceId,
                               __tgt_async_info *AsyncInfoPtr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   assert(AsyncInfoPtr && "async_info_ptr is nullptr");
   assert(AsyncInfoPtr->Queue && "async_info_ptr->Queue is nullptr");
   // NOTE: We don't need to set context for stream sync.
-  return DeviceRTL.synchronize(DeviceId, AsyncInfoPtr);
+  return DeviceRTL->synchronize(DeviceId, AsyncInfoPtr);
 }
 
 void __tgt_rtl_set_info_flag(uint32_t NewInfoLevel) {
   std::atomic<uint32_t> &InfoLevel = getInfoLevelInternal();
   InfoLevel.store(NewInfoLevel);
 }
 
 void __tgt_rtl_print_device_info(int32_t DeviceId) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   // NOTE: We don't need to set context for print device info.
-  DeviceRTL.printDeviceInfo(DeviceId);
+  DeviceRTL->printDeviceInfo(DeviceId);
 }
 
 int32_t __tgt_rtl_create_event(int32_t DeviceId, void **Event) {
   assert(Event && "event is nullptr");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.createEvent(DeviceId, Event);
+  return DeviceRTL->createEvent(DeviceId, Event);
 }
 
 int32_t __tgt_rtl_record_event(int32_t DeviceId, void *EventPtr,
                                __tgt_async_info *AsyncInfoPtr) {
   assert(AsyncInfoPtr && "async_info_ptr is nullptr");
   assert(AsyncInfoPtr->Queue && "async_info_ptr->Queue is nullptr");
   assert(EventPtr && "event_ptr is nullptr");
   // NOTE: We might not need to set context for event record.
   return recordEvent(EventPtr, AsyncInfoPtr);
 }
 
 int32_t __tgt_rtl_wait_event(int32_t DeviceId, void *EventPtr,
                              __tgt_async_info *AsyncInfoPtr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   assert(AsyncInfoPtr && "async_info_ptr is nullptr");
   assert(EventPtr && "event is nullptr");
   // If we don't have a queue we need to set the context.
-  if (!AsyncInfoPtr->Queue && DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (!AsyncInfoPtr->Queue &&
+      DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
-  return DeviceRTL.waitEvent(DeviceId, AsyncInfoPtr, EventPtr);
+  return DeviceRTL->waitEvent(DeviceId, AsyncInfoPtr, EventPtr);
 }
 
 int32_t __tgt_rtl_sync_event(int32_t DeviceId, void *EventPtr) {
   assert(EventPtr && "event is nullptr");
   // NOTE: We might not need to set context for event sync.
   return syncEvent(EventPtr);
 }
 
 int32_t __tgt_rtl_destroy_event(int32_t DeviceId, void *EventPtr) {
   assert(EventPtr && "event is nullptr");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.destroyEvent(DeviceId, EventPtr);
+  return DeviceRTL->destroyEvent(DeviceId, EventPtr);
 }
 
 int32_t __tgt_rtl_release_async_info(int32_t DeviceId,
                                      __tgt_async_info *AsyncInfo) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   assert(AsyncInfo && "async_info is nullptr");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.releaseAsyncInfo(DeviceId, AsyncInfo);
+  return DeviceRTL->releaseAsyncInfo(DeviceId, AsyncInfo);
 }
 
 int32_t __tgt_rtl_init_async_info(int32_t DeviceId,
                                   __tgt_async_info **AsyncInfo) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   assert(AsyncInfo && "async_info is nullptr");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.initAsyncInfo(DeviceId, AsyncInfo);
+  return DeviceRTL->initAsyncInfo(DeviceId, AsyncInfo);
 }
 
 int32_t __tgt_rtl_init_device_info(int32_t DeviceId,
                                    __tgt_device_info *DeviceInfoPtr,
                                    const char **ErrStr) {
-  assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceRTL->isValidDeviceId(DeviceId) && "device_id is invalid");
   assert(DeviceInfoPtr && "device_info_ptr is nullptr");
 
-  if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
+  if (DeviceRTL->setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.initDeviceInfo(DeviceId, DeviceInfoPtr, ErrStr);
+  return DeviceRTL->initDeviceInfo(DeviceId, DeviceInfoPtr, ErrStr);
 }
 
 #ifdef __cplusplus
 }
 #endif

openmp/libomptarget/plugins/generic-elf-64bit/src/rtl.cpp

 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <dlfcn.h>
 #include <ffi.h>
 #include <gelf.h>
 #include <link.h>
 #include <list>
+#include <memory>
 #include <string>
 #include <vector>
 
 #include "Debug.h"
 #include "omptargetplugin.h"
 
 #ifndef TARGET_NAME
 #define TARGET_NAME Generic ELF - 64bit
for (auto &Lib : DynLibs) {
       if (Lib.Handle) {
         dlclose(Lib.Handle);
         remove(Lib.FileName.c_str());
       }
     }
   }
 };
 
-static RTLDeviceInfoTy DeviceInfo(NUMBER_OF_DEVICES);
+RTLDeviceInfoTy *DeviceInfo = nullptr;
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
+int32_t __tgt_rtl_init_plugin() {
+  DeviceInfo = new RTLDeviceInfoTy(NUMBER_OF_DEVICES);
+  return DeviceInfo ? OFFLOAD_SUCCESS : OFFLOAD_FAIL;
+}
+
+int32_t __tgt_rtl_deinit_plugin() {
+  if (DeviceInfo)
+    delete DeviceInfo;
+  return OFFLOAD_SUCCESS;
+}
+
 int32_t __tgt_rtl_is_valid_binary(__tgt_device_image *Image) {
 // If we don't have a valid ELF ID we can just fail.
 #if TARGET_ELF_ID < 1
   return 0;
 #else
   return elf_check_machine(Image, TARGET_ELF_ID);
 #endif
 }
__tgt_target_table *__tgt_rtl_load_binary(int32_t DeviceId,
   DynLibTy Lib = {TmpName, dlopen(TmpName, RTLD_LAZY)};
 
   if (!Lib.Handle) {
     DP("Target library loading error: %s\n", dlerror());
     elf_end(E);
     return NULL;
   }
 
-  DeviceInfo.DynLibs.push_back(Lib);
+  DeviceInfo->DynLibs.push_back(Lib);
 
   struct link_map *LibInfo = (struct link_map *)Lib.Handle;
 
   // The place where the entries info is loaded is the library base address
   // plus the offset determined from the ELF file.
   Elf64_Addr EntriesAddr = LibInfo->l_addr + EntriesOffset;
 
   DP("Pointer to first entry to be loaded is (" DPxMOD ").\n",
      DPxPTR(EntriesAddr));
 
   // Table of pointers to all the entries in the target.
   __tgt_offload_entry *EntriesTable = (__tgt_offload_entry *)EntriesAddr;
 
   __tgt_offload_entry *EntriesBegin = &EntriesTable[0];
   __tgt_offload_entry *EntriesEnd = EntriesBegin + NumEntries;
 
   if (!EntriesBegin) {
     DP("Can't obtain entries begin\n");
     elf_end(E);
     return NULL;
   }
 
   DP("Entries table range is (" DPxMOD ")->(" DPxMOD ")\n",
      DPxPTR(EntriesBegin), DPxPTR(EntriesEnd));
-  DeviceInfo.createOffloadTable(DeviceId, EntriesBegin, EntriesEnd);
+  DeviceInfo->createOffloadTable(DeviceId, EntriesBegin, EntriesEnd);
 
   elf_end(E);
 
-  return DeviceInfo.getOffloadEntriesTable(DeviceId);
+  return DeviceInfo->getOffloadEntriesTable(DeviceId);
 }
 
 void __tgt_rtl_print_device_info(int32_t DeviceId) {
   printf("    This is a generic-elf-64bit device\n");
 }
 
 // Sample implementation of explicit memory allocator. For this plugin all kinds
 // are equivalent to each other.

openmp/libomptarget/test/offloading/global_constructor.cpp

 // RUN: %libomptarget-compilexx-generic && %libomptarget-run-generic | %fcheck-generic
 
 #include <cstdio>
 
 int foo() { return 1; }
 
 class C {
 public:
   C() : x(foo()) {}
+  ~C() { foo(); }
 
   int x;
 };
 
 C c;
 #pragma omp declare target(c)
 
 int main() {
   int x = 0;
 #pragma omp target map(from : x)
   { x = c.x; }
 
   // CHECK: PASS
   if (x == 1)
     printf("PASS\n");
 }