diff --git a/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp b/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
--- a/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
@@ -24,18 +24,6 @@
   return core::Runtime::GetMachineInfo();
 }
 
-/*
- * Modules
- */
-atmi_status_t atmi_module_register_from_memory_to_place(
-    void *module_bytes, size_t module_size, atmi_place_t place,
-    atmi_status_t (*on_deserialized_data)(void *data, size_t size,
-                                          void *cb_state),
-    void *cb_state) {
-  return core::Runtime::getInstance().RegisterModuleFromMemory(
-      module_bytes, module_size, place, on_deserialized_data, cb_state);
-}
-
 /*
  * Data
  */
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/rt.h b/openmp/libomptarget/plugins/amdgpu/impl/rt.h
--- a/openmp/libomptarget/plugins/amdgpu/impl/rt.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/rt.h
@@ -59,7 +59,7 @@
       void *, size_t, atmi_place_t,
       atmi_status_t (*on_deserialized_data)(void *data, size_t size,
                                             void *cb_state),
-      void *cb_state);
+      void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
 
   // data
   static atmi_status_t Memcpy(hsa_signal_t, void *, const void *, size_t);
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
--- a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
@@ -132,8 +132,6 @@
 
 std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
 
-static std::vector<hsa_executable_t> g_executables;
-
 std::map<std::string, std::string> KernelNameMap;
 std::vector<std::map<std::string, atl_kernel_info_t>> KernelInfoTable;
 std::vector<std::map<std::string, atl_symbol_info_t>> SymbolInfoTable;
@@ -204,15 +202,6 @@
 atmi_status_t Runtime::Finalize() {
   hsa_status_t err;
 
-  for (uint32_t i = 0; i < g_executables.size(); i++) {
-    err = hsa_executable_destroy(g_executables[i]);
-    if (err != HSA_STATUS_SUCCESS) {
-      printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
-             "Destroying executable", get_error_string(err));
-      exit(1);
-    }
-  }
-
   for (uint32_t i = 0; i < SymbolInfoTable.size(); i++) {
     SymbolInfoTable[i].clear();
   }
@@ -1170,7 +1159,7 @@
     void *module_bytes, size_t module_size, atmi_place_t place,
     atmi_status_t (*on_deserialized_data)(void *data, size_t size,
                                           void *cb_state),
-    void *cb_state) {
+    void *cb_state, std::vector<hsa_executable_t> &HSAExecutables) {
   hsa_status_t err;
   int gpu = place.device_id;
   assert(gpu >= 0);
@@ -1269,7 +1258,7 @@
     }
 
     // save the executable and destroy during finalize
-    g_executables.push_back(executable);
+    HSAExecutables.push_back(executable);
     return ATMI_STATUS_SUCCESS;
   } else {
     return ATMI_STATUS_ERROR;
diff --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -328,6 +328,8 @@
   // Resource pools
   SignalPoolT FreeSignalPool;
 
+  std::vector<hsa_executable_t> HSAExecutables;
+
   struct atmiFreePtrDeletor {
     void operator()(void *p) {
       atmi_free(p); // ignore failure to free
@@ -538,6 +540,18 @@
     RequiresFlags = OMP_REQ_UNDEFINED;
   }
 
+  void DestroyHSAExecutables() {
+    hsa_status_t Err;
+    for (uint32_t I = 0; I < HSAExecutables.size(); I++) {
+      Err = hsa_executable_destroy(HSAExecutables[I]);
+      if (Err != HSA_STATUS_SUCCESS) {
+        printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+               "Destroying executable", get_error_string(Err));
+        return;
+      }
+    }
+  }
+
   ~RTLDeviceInfoTy() {
     DP("Finalizing the HSA-ATMI DeviceInfo.\n");
     // Run destructors on types that use HSA before
@@ -546,6 +560,8 @@
     KernelArgPoolMap.clear();
     // Terminate hostrpc before finalizing ATMI
     hostrpc_terminate();
+
+    DestroyHSAExecutables();
     atmi_finalize();
   }
 };
@@ -971,15 +987,16 @@
 }
 
 template <typename C>
-atmi_status_t module_register_from_memory_to_place(void *module_bytes,
-                                                   size_t module_size,
-                                                   atmi_place_t place, C cb) {
+atmi_status_t module_register_from_memory_to_place(
+    void *module_bytes, size_t module_size, atmi_place_t place, C cb,
+    std::vector<hsa_executable_t> &HSAExecutables) {
   auto L = [](void *data, size_t size, void *cb_state) -> atmi_status_t {
     C *unwrapped = static_cast<C *>(cb_state);
     return (*unwrapped)(data, size);
   };
-  return atmi_module_register_from_memory_to_place(
-      module_bytes, module_size, place, L, static_cast<void *>(&cb));
+  return core::Runtime::RegisterModuleFromMemory(
+      module_bytes, module_size, place, L, static_cast<void *>(&cb),
+      HSAExecutables);
 }
 } // namespace
 
@@ -1180,9 +1197,8 @@
 
     atmi_status_t err = module_register_from_memory_to_place(
         (void *)image->ImageStart, img_size, get_gpu_place(device_id),
-        [&](void *data, size_t size) {
-          return env.before_loading(data, size);
-        });
+        [&](void *data, size_t size) { return env.before_loading(data, size); },
+        DeviceInfo.HSAExecutables);
 
     check("Module registering", err);
     if (err != ATMI_STATUS_SUCCESS) {