[LIBOMPTARGET][WIP]Fix PR45139: Bug in mixing Python and OpenMP target offload.

[ABataev]

Summary

Dynamically initialize objects in cuda/generic plugins to avoid issue
with not initialized data upob libraries loading.

[jdoerfert]

LGTM.

[JonChesterfield]

This feels familiar.

It's a shame to be heap allocating all this stuff in order to keep initialisation working. Can we lazily init instead, in singleton fashion?

[ABataev]

In D75926#1914977, @jdoerfert wrote:

LGTM.

It does not work, actually. Still investigating

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

for (auto &ctx : Contexts)
         if (err != CUDA_SUCCESS) {
           DP("Error when destroying CUDA context\n");
           CUDA_ERR_STRING(err);
         }
       }
   }
 };
 
-static RTLDeviceInfoTy DeviceInfo;
+static RTLDeviceInfoTy *DeviceInfo;
+
+__attribute__((constructor(101))) void init() {
+  DP("Init CUDA plugin!\n");
+  DeviceInfo = new RTLDeviceInfoTy();
+}
+
+__attribute__((destructor(101))) void deinit() {
+  DP("Deinit CUDA plugin!\n");
+  delete DeviceInfo;
+}
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 int32_t __tgt_rtl_is_valid_binary(__tgt_device_image *image) {
   return elf_check_machine(image, 190); // EM_CUDA = 190.
 }
 
-int32_t __tgt_rtl_number_of_devices() { return DeviceInfo.NumberOfDevices; }
+int32_t __tgt_rtl_number_of_devices() { return DeviceInfo->NumberOfDevices; }
 
 int64_t __tgt_rtl_init_requires(int64_t RequiresFlags) {
   DP("Init requires flags to %ld\n", RequiresFlags);
-  DeviceInfo.RequiresFlags = RequiresFlags;
+  DeviceInfo->RequiresFlags = RequiresFlags;
   return RequiresFlags;
 }
 
 int32_t __tgt_rtl_init_device(int32_t device_id) {
 
   CUdevice cuDevice;
   DP("Getting device %d\n", device_id);
   CUresult err = cuDeviceGet(&cuDevice, device_id);
   if (err != CUDA_SUCCESS) {
     DP("Error when getting CUDA device with id = %d\n", device_id);
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
   // Create the context and save it to use whenever this device is selected.
-  err = cuCtxCreate(&DeviceInfo.Contexts[device_id], CU_CTX_SCHED_BLOCKING_SYNC,
+  err = cuCtxCreate(&DeviceInfo->Contexts[device_id], CU_CTX_SCHED_BLOCKING_SYNC,
                     cuDevice);
   if (err != CUDA_SUCCESS) {
     DP("Error when creating a CUDA context\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
-  err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  err = cuCtxSetCurrent(DeviceInfo->Contexts[device_id]);
   if (err != CUDA_SUCCESS) {
     DP("Error when setting current CUDA context\n");
     CUDA_ERR_STRING(err);
   }
 
-  for (CUstream &Stream : DeviceInfo.Streams[device_id]) {
+  for (CUstream &Stream : DeviceInfo->Streams[device_id]) {
     err = cuStreamCreate(&Stream, CU_STREAM_NON_BLOCKING);
     if (err != CUDA_SUCCESS) {
       DP("Error when creating CUDA stream\n");
       CUDA_ERR_STRING(err);
     }
   }
 
   // Query attributes to determine number of threads/block and blocks/grid.
   int maxGridDimX;
   err = cuDeviceGetAttribute(&maxGridDimX, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X,
                              cuDevice);
   if (err != CUDA_SUCCESS) {
     DP("Error getting max grid dimension, use default\n");
-    DeviceInfo.BlocksPerGrid[device_id] = RTLDeviceInfoTy::DefaultNumTeams;
+    DeviceInfo->BlocksPerGrid[device_id] = RTLDeviceInfoTy::DefaultNumTeams;
   } else if (maxGridDimX <= RTLDeviceInfoTy::HardTeamLimit) {
-    DeviceInfo.BlocksPerGrid[device_id] = maxGridDimX;
+    DeviceInfo->BlocksPerGrid[device_id] = maxGridDimX;
     DP("Using %d CUDA blocks per grid\n", maxGridDimX);
   } else {
-    DeviceInfo.BlocksPerGrid[device_id] = RTLDeviceInfoTy::HardTeamLimit;
+    DeviceInfo->BlocksPerGrid[device_id] = RTLDeviceInfoTy::HardTeamLimit;
     DP("Max CUDA blocks per grid %d exceeds the hard team limit %d, capping "
        "at the hard limit\n",
        maxGridDimX, RTLDeviceInfoTy::HardTeamLimit);
   }
 
   // We are only exploiting threads along the x axis.
   int maxBlockDimX;
   err = cuDeviceGetAttribute(&maxBlockDimX, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X,
                              cuDevice);
   if (err != CUDA_SUCCESS) {
     DP("Error getting max block dimension, use default\n");
-    DeviceInfo.ThreadsPerBlock[device_id] = RTLDeviceInfoTy::DefaultNumThreads;
+    DeviceInfo->ThreadsPerBlock[device_id] = RTLDeviceInfoTy::DefaultNumThreads;
   } else if (maxBlockDimX <= RTLDeviceInfoTy::HardThreadLimit) {
-    DeviceInfo.ThreadsPerBlock[device_id] = maxBlockDimX;
+    DeviceInfo->ThreadsPerBlock[device_id] = maxBlockDimX;
     DP("Using %d CUDA threads per block\n", maxBlockDimX);
   } else {
-    DeviceInfo.ThreadsPerBlock[device_id] = RTLDeviceInfoTy::HardThreadLimit;
+    DeviceInfo->ThreadsPerBlock[device_id] = RTLDeviceInfoTy::HardThreadLimit;
     DP("Max CUDA threads per block %d exceeds the hard thread limit %d, capping"
        "at the hard limit\n",
        maxBlockDimX, RTLDeviceInfoTy::HardThreadLimit);
   }
 
   int warpSize;
   err =
       cuDeviceGetAttribute(&warpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, cuDevice);
   if (err != CUDA_SUCCESS) {
     DP("Error getting warp size, assume default\n");
-    DeviceInfo.WarpSize[device_id] = 32;
+    DeviceInfo->WarpSize[device_id] = 32;
   } else {
-    DeviceInfo.WarpSize[device_id] = warpSize;
+    DeviceInfo->WarpSize[device_id] = warpSize;
   }
 
   // Adjust teams to the env variables
-  if (DeviceInfo.EnvTeamLimit > 0 &&
-      DeviceInfo.BlocksPerGrid[device_id] > DeviceInfo.EnvTeamLimit) {
-    DeviceInfo.BlocksPerGrid[device_id] = DeviceInfo.EnvTeamLimit;
+  if (DeviceInfo->EnvTeamLimit > 0 &&
+      DeviceInfo->BlocksPerGrid[device_id] > DeviceInfo->EnvTeamLimit) {
+    DeviceInfo->BlocksPerGrid[device_id] = DeviceInfo->EnvTeamLimit;
     DP("Capping max CUDA blocks per grid to OMP_TEAM_LIMIT=%d\n",
-        DeviceInfo.EnvTeamLimit);
+        DeviceInfo->EnvTeamLimit);
   }
 
   DP("Max number of CUDA blocks %d, threads %d & warp size %d\n",
-     DeviceInfo.BlocksPerGrid[device_id], DeviceInfo.ThreadsPerBlock[device_id],
-     DeviceInfo.WarpSize[device_id]);
+     DeviceInfo->BlocksPerGrid[device_id],
+     DeviceInfo->ThreadsPerBlock[device_id], DeviceInfo->WarpSize[device_id]);
 
   // Set default number of teams
-  if (DeviceInfo.EnvNumTeams > 0) {
-    DeviceInfo.NumTeams[device_id] = DeviceInfo.EnvNumTeams;
+  if (DeviceInfo->EnvNumTeams > 0) {
+    DeviceInfo->NumTeams[device_id] = DeviceInfo->EnvNumTeams;
     DP("Default number of teams set according to environment %d\n",
-        DeviceInfo.EnvNumTeams);
+        DeviceInfo->EnvNumTeams);
   } else {
-    DeviceInfo.NumTeams[device_id] = RTLDeviceInfoTy::DefaultNumTeams;
+    DeviceInfo->NumTeams[device_id] = RTLDeviceInfoTy::DefaultNumTeams;
     DP("Default number of teams set according to library's default %d\n",
         RTLDeviceInfoTy::DefaultNumTeams);
   }
-  if (DeviceInfo.NumTeams[device_id] > DeviceInfo.BlocksPerGrid[device_id]) {
-    DeviceInfo.NumTeams[device_id] = DeviceInfo.BlocksPerGrid[device_id];
+  if (DeviceInfo->NumTeams[device_id] > DeviceInfo->BlocksPerGrid[device_id]) {
+    DeviceInfo->NumTeams[device_id] = DeviceInfo->BlocksPerGrid[device_id];
     DP("Default number of teams exceeds device limit, capping at %d\n",
-        DeviceInfo.BlocksPerGrid[device_id]);
+        DeviceInfo->BlocksPerGrid[device_id]);
   }
 
   // Set default number of threads
-  DeviceInfo.NumThreads[device_id] = RTLDeviceInfoTy::DefaultNumThreads;
+  DeviceInfo->NumThreads[device_id] = RTLDeviceInfoTy::DefaultNumThreads;
   DP("Default number of threads set according to library's default %d\n",
           RTLDeviceInfoTy::DefaultNumThreads);
-  if (DeviceInfo.NumThreads[device_id] >
-      DeviceInfo.ThreadsPerBlock[device_id]) {
-    DeviceInfo.NumTeams[device_id] = DeviceInfo.ThreadsPerBlock[device_id];
+  if (DeviceInfo->NumThreads[device_id] >
+      DeviceInfo->ThreadsPerBlock[device_id]) {
+    DeviceInfo->NumTeams[device_id] = DeviceInfo->ThreadsPerBlock[device_id];
     DP("Default number of threads exceeds device limit, capping at %d\n",
-        DeviceInfo.ThreadsPerBlock[device_id]);
+        DeviceInfo->ThreadsPerBlock[device_id]);
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 __tgt_target_table *__tgt_rtl_load_binary(int32_t device_id,
     __tgt_device_image *image) {
 
   // Set the context we are using.
-  CUresult err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  CUresult err = cuCtxSetCurrent(DeviceInfo->Contexts[device_id]);
   if (err != CUDA_SUCCESS) {
     DP("Error when setting a CUDA context for device %d\n", device_id);
     CUDA_ERR_STRING(err);
     return NULL;
   }
 
   // Clear the offload table as we are going to create a new one.
-  DeviceInfo.clearOffloadEntriesTable(device_id);
+  DeviceInfo->clearOffloadEntriesTable(device_id);
 
   // Create the module and extract the function pointers.
 
   CUmodule cumod;
   DP("Load data from image " DPxMOD "\n", DPxPTR(image->ImageStart));
   err = cuModuleLoadDataEx(&cumod, image->ImageStart, 0, NULL, NULL);
   if (err != CUDA_SUCCESS) {
     DP("Error when loading CUDA module\n");
     CUDA_ERR_STRING(err);
     return NULL;
   }
 
   DP("CUDA module successfully loaded!\n");
-  DeviceInfo.Modules.push_back(cumod);
+  DeviceInfo->Modules.push_back(cumod);
 
   // Find the symbols in the module by name.
   __tgt_offload_entry *HostBegin = image->EntriesBegin;
   __tgt_offload_entry *HostEnd = image->EntriesEnd;
 
   for (__tgt_offload_entry *e = HostBegin; e != HostEnd; ++e) {
 
     if (!e->addr) {
if (e->size) {
       // memory is activated via the requires directive, the variable
       // can be used directly from the host in both cases.
       // TODO: when variables types other than to or link are added,
       // the below condition should be changed to explicitly
       // check for to and link variables types:
       //  (DeviceInfo.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
       //   (e->flags & OMP_DECLARE_TARGET_LINK ||
       //    e->flags == OMP_DECLARE_TARGET_TO))
-      if (DeviceInfo.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
+      if (DeviceInfo->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
         // If unified memory is present any target link or to variables
         // can access host addresses directly. There is no longer a
         // need for device copies.
         cuMemcpyHtoD(cuptr, e->addr, sizeof(void *));
         DP("Copy linked variable host address (" DPxMOD ")"
            "to device address (" DPxMOD ")\n",
           DPxPTR(*((void**)e->addr)), DPxPTR(cuptr));
       }
 
-      DeviceInfo.addOffloadEntry(device_id, entry);
+      DeviceInfo->addOffloadEntry(device_id, entry);
 
       continue;
     }
 
     CUfunction fun;
     err = cuModuleGetFunction(&fun, cumod, e->name);
 
     if (err != CUDA_SUCCESS) {
if (err == CUDA_SUCCESS) {
           "GENERIC (1)\n", ExecModeName);
       CUDA_ERR_STRING(err);
     }
 
     KernelsList.push_back(KernelTy(fun, ExecModeVal));
 
     __tgt_offload_entry entry = *e;
     entry.addr = (void *)&KernelsList.back();
-    DeviceInfo.addOffloadEntry(device_id, entry);
+    DeviceInfo->addOffloadEntry(device_id, entry);
   }
 
   // send device environment data to the device
   {
     omptarget_device_environmentTy device_env;
 
     device_env.debug_level = 0;
 
if (err == CUDA_SUCCESS) {
 
       DP("Sending global device environment data %zu bytes\n", (size_t)cusize);
     } else {
       DP("Finding global device environment '%s' - symbol missing.\n", device_env_Name);
       DP("Continue, considering this is a device RTL which does not accept environment setting.\n");
     }
   }
 
-  return DeviceInfo.getOffloadEntriesTable(device_id);
+  return DeviceInfo->getOffloadEntriesTable(device_id);
 }
 
 void *__tgt_rtl_data_alloc(int32_t device_id, int64_t size, void *hst_ptr) {
   if (size == 0) {
     return NULL;
   }
 
   // Set the context we are using.
-  CUresult err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  CUresult err = cuCtxSetCurrent(DeviceInfo->Contexts[device_id]);
   if (err != CUDA_SUCCESS) {
     DP("Error while trying to set CUDA current context\n");
     CUDA_ERR_STRING(err);
     return NULL;
   }
 
   CUdeviceptr ptr;
   err = cuMemAlloc(&ptr, size);
   if (err != CUDA_SUCCESS) {
     DP("Error while trying to allocate %d\n", err);
     CUDA_ERR_STRING(err);
     return NULL;
   }
 
   void *vptr = (void *)ptr;
   return vptr;
 }
 
 int32_t __tgt_rtl_data_submit(int32_t device_id, void *tgt_ptr, void *hst_ptr,
     int64_t size) {
   // Set the context we are using.
-  CUresult err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  CUresult err = cuCtxSetCurrent(DeviceInfo->Contexts[device_id]);
   if (err != CUDA_SUCCESS) {
     DP("Error when setting CUDA context\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
-  CUstream &Stream = DeviceInfo.getNextStream(device_id);
+  CUstream &Stream = DeviceInfo->getNextStream(device_id);
 
   err = cuMemcpyHtoDAsync((CUdeviceptr)tgt_ptr, hst_ptr, size, Stream);
   if (err != CUDA_SUCCESS) {
     DP("Error when copying data from host to device. Pointers: host = " DPxMOD
        ", device = " DPxMOD ", size = %" PRId64 "\n",
        DPxPTR(hst_ptr), DPxPTR(tgt_ptr), size);
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
int32_t __tgt_rtl_data_submit(int32_t device_id, void *tgt_ptr, void *hst_ptr,
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 int32_t __tgt_rtl_data_retrieve(int32_t device_id, void *hst_ptr, void *tgt_ptr,
     int64_t size) {
   // Set the context we are using.
-  CUresult err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  CUresult err = cuCtxSetCurrent(DeviceInfo->Contexts[device_id]);
   if (err != CUDA_SUCCESS) {
     DP("Error when setting CUDA context\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
-  CUstream &Stream = DeviceInfo.getNextStream(device_id);
+  CUstream &Stream = DeviceInfo->getNextStream(device_id);
 
   err = cuMemcpyDtoHAsync(hst_ptr, (CUdeviceptr)tgt_ptr, size, Stream);
   if (err != CUDA_SUCCESS) {
     DP("Error when copying data from device to host. Pointers: host = " DPxMOD
        ", device = " DPxMOD ", size = %" PRId64 "\n",
        DPxPTR(hst_ptr), DPxPTR(tgt_ptr), size);
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
   err = cuStreamSynchronize(Stream);
   if (err != CUDA_SUCCESS) {
     DP("Error when synchronizing async data transfer from device to host. "
        "Pointers: host = " DPxMOD ", device = " DPxMOD ", size = %" PRId64 "\n",
        DPxPTR(hst_ptr), DPxPTR(tgt_ptr), size);
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 int32_t __tgt_rtl_data_delete(int32_t device_id, void *tgt_ptr) {
   // Set the context we are using.
-  CUresult err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  CUresult err = cuCtxSetCurrent(DeviceInfo->Contexts[device_id]);
   if (err != CUDA_SUCCESS) {
     DP("Error when setting CUDA context\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
   err = cuMemFree((CUdeviceptr)tgt_ptr);
   if (err != CUDA_SUCCESS) {
     DP("Error when freeing CUDA memory\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
   return OFFLOAD_SUCCESS;
 }
 
 int32_t __tgt_rtl_run_target_team_region(int32_t device_id, void *tgt_entry_ptr,
     void **tgt_args, ptrdiff_t *tgt_offsets, int32_t arg_num, int32_t team_num,
     int32_t thread_limit, uint64_t loop_tripcount) {
   // Set the context we are using.
-  CUresult err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  CUresult err = cuCtxSetCurrent(DeviceInfo->Contexts[device_id]);
   if (err != CUDA_SUCCESS) {
     DP("Error when setting CUDA context\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
   // All args are references.
   std::vector<void *> args(arg_num);
   std::vector<void *> ptrs(arg_num);
 
   for (int32_t i = 0; i < arg_num; ++i) {
     ptrs[i] = (void *)((intptr_t)tgt_args[i] + tgt_offsets[i]);
     args[i] = &ptrs[i];
   }
 
   KernelTy *KernelInfo = (KernelTy *)tgt_entry_ptr;
 
   int cudaThreadsPerBlock;
 
   if (thread_limit > 0) {
     cudaThreadsPerBlock = thread_limit;
     DP("Setting CUDA threads per block to requested %d\n", thread_limit);
     // Add master warp if necessary
     if (KernelInfo->ExecutionMode == GENERIC) {
-      cudaThreadsPerBlock += DeviceInfo.WarpSize[device_id];
-      DP("Adding master warp: +%d threads\n", DeviceInfo.WarpSize[device_id]);
+      cudaThreadsPerBlock += DeviceInfo->WarpSize[device_id];
+      DP("Adding master warp: +%d threads\n", DeviceInfo->WarpSize[device_id]);
     }
   } else {
-    cudaThreadsPerBlock = DeviceInfo.NumThreads[device_id];
+    cudaThreadsPerBlock = DeviceInfo->NumThreads[device_id];
     DP("Setting CUDA threads per block to default %d\n",
-        DeviceInfo.NumThreads[device_id]);
+        DeviceInfo->NumThreads[device_id]);
   }
 
-  if (cudaThreadsPerBlock > DeviceInfo.ThreadsPerBlock[device_id]) {
-    cudaThreadsPerBlock = DeviceInfo.ThreadsPerBlock[device_id];
+  if (cudaThreadsPerBlock > DeviceInfo->ThreadsPerBlock[device_id]) {
+    cudaThreadsPerBlock = DeviceInfo->ThreadsPerBlock[device_id];
     DP("Threads per block capped at device limit %d\n",
-        DeviceInfo.ThreadsPerBlock[device_id]);
+        DeviceInfo->ThreadsPerBlock[device_id]);
   }
 
   int kernel_limit;
   err = cuFuncGetAttribute(&kernel_limit,
       CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, KernelInfo->Func);
   if (err == CUDA_SUCCESS) {
     if (kernel_limit < cudaThreadsPerBlock) {
       cudaThreadsPerBlock = kernel_limit;
       DP("Threads per block capped at kernel limit %d\n", kernel_limit);
     }
   }
 
   int cudaBlocksPerGrid;
   if (team_num <= 0) {
-    if (loop_tripcount > 0 && DeviceInfo.EnvNumTeams < 0) {
+    if (loop_tripcount > 0 && DeviceInfo->EnvNumTeams < 0) {
       if (KernelInfo->ExecutionMode == SPMD) {
         // We have a combined construct, i.e. `target teams distribute parallel
         // for [simd]`. We launch so many teams so that each thread will
         // execute one iteration of the loop.
         // round up to the nearest integer
         cudaBlocksPerGrid = ((loop_tripcount - 1) / cudaThreadsPerBlock) + 1;
       } else {
         // If we reach this point, then we have a non-combined construct, i.e.
if (loop_tripcount > 0 && DeviceInfo->EnvNumTeams < 0) {
         // Threads within a team will execute the iterations of the `parallel`
         // loop.
         cudaBlocksPerGrid = loop_tripcount;
       }
       DP("Using %d teams due to loop trip count %" PRIu64 " and number of "
           "threads per block %d\n", cudaBlocksPerGrid, loop_tripcount,
           cudaThreadsPerBlock);
     } else {
-      cudaBlocksPerGrid = DeviceInfo.NumTeams[device_id];
-      DP("Using default number of teams %d\n", DeviceInfo.NumTeams[device_id]);
+      cudaBlocksPerGrid = DeviceInfo->NumTeams[device_id];
+      DP("Using default number of teams %d\n", DeviceInfo->NumTeams[device_id]);
     }
-  } else if (team_num > DeviceInfo.BlocksPerGrid[device_id]) {
-    cudaBlocksPerGrid = DeviceInfo.BlocksPerGrid[device_id];
+  } else if (team_num > DeviceInfo->BlocksPerGrid[device_id]) {
+    cudaBlocksPerGrid = DeviceInfo->BlocksPerGrid[device_id];
     DP("Capping number of teams to team limit %d\n",
-        DeviceInfo.BlocksPerGrid[device_id]);
+        DeviceInfo->BlocksPerGrid[device_id]);
   } else {
     cudaBlocksPerGrid = team_num;
     DP("Using requested number of teams %d\n", team_num);
   }
 
   // Run on the device.
   DP("Launch kernel with %d blocks and %d threads\n", cudaBlocksPerGrid,
      cudaThreadsPerBlock);
 
-  CUstream &Stream = DeviceInfo.getNextStream(device_id);
+  CUstream &Stream = DeviceInfo->getNextStream(device_id);
 
   err = cuLaunchKernel(KernelInfo->Func, cudaBlocksPerGrid, 1, 1,
                        cudaThreadsPerBlock, 1, 1, 0 /*bytes of shared memory*/,
                        Stream, &args[0], 0);
   if (err != CUDA_SUCCESS) {
     DP("Device kernel launch failed!\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;

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

for (auto &lib : DynLibs) {
       if (lib.Handle) {
         dlclose(lib.Handle);
         remove(lib.FileName);
       }
     }
   }
 };
 
-static RTLDeviceInfoTy DeviceInfo(NUMBER_OF_DEVICES);
+static RTLDeviceInfoTy *DeviceInfo;
+
+__attribute__((constructor(101))) void init() {
+  DP("Init generic plugin!\n");
+  DeviceInfo = new RTLDeviceInfoTy(NUMBER_OF_DEVICES);
+}
+
+__attribute__((destructor(101))) void deinit() {
+  DP("Deinit generic plugin!\n");
+  delete DeviceInfo;
+}
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 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
__tgt_target_table *__tgt_rtl_load_binary(int32_t device_id,
   DynLibTy Lib = {tmp_name, dlopen(tmp_name, 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 entries_addr = libInfo->l_addr + entries_offset;
 
   DP("Pointer to first entry to be loaded is (" DPxMOD ").\n",
       DPxPTR(entries_addr));
 
   // Table of pointers to all the entries in the target.
   __tgt_offload_entry *entries_table = (__tgt_offload_entry *)entries_addr;
 
   __tgt_offload_entry *entries_begin = &entries_table[0];
   __tgt_offload_entry *entries_end = entries_begin + NumEntries;
 
   if (!entries_begin) {
     DP("Can't obtain entries begin\n");
     elf_end(e);
     return NULL;
   }
 
   DP("Entries table range is (" DPxMOD ")->(" DPxMOD ")\n",
       DPxPTR(entries_begin), DPxPTR(entries_end));
-  DeviceInfo.createOffloadTable(device_id, entries_begin, entries_end);
+  DeviceInfo->createOffloadTable(device_id, entries_begin, entries_end);
 
   elf_end(e);
 
-  return DeviceInfo.getOffloadEntriesTable(device_id);
+  return DeviceInfo->getOffloadEntriesTable(device_id);
 }
 
 void *__tgt_rtl_data_alloc(int32_t device_id, int64_t size, void *hst_ptr) {
   void *ptr = malloc(size);
   return ptr;
 }
 
 int32_t __tgt_rtl_data_submit(int32_t device_id, void *tgt_ptr, void *hst_ptr,