[AMDGPU][Libomptarget] Rename & move g_executables to private

This patch moves g_executables to private member of Runtime class
and is renamed to HSAExecutables following LLVM naming convention.

This movement required making Runtime::Initialize and Runtime::Finalize
non-static. Verified the correctness of this change by running
libomptarget tests on gfx906.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D102600

openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp

@@ -24,18 +24,6 @@ atmi_machine_t *atmi_machine_get_info() {
   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
  */

openmp/libomptarget/plugins/amdgpu/impl/rt.h

@@ -59,7 +59,7 @@ public:
       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);

openmp/libomptarget/plugins/amdgpu/impl/system.cpp

@@ -132,8 +132,6 @@ ATLMachine g_atl_machine;
 
 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::Initialize() {
 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 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
     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 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
     }
 
     // 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;

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

@@ -328,6 +328,8 @@ public:
   // 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 @@ public:
     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 @@ public:
     KernelArgPoolMap.clear();
     // Terminate hostrpc before finalizing ATMI
     hostrpc_terminate();
+
+    DestroyHSAExecutables();
     atmi_finalize();
   }
 };
@@ -971,15 +987,16 @@ atmi_status_t interop_get_symbol_info(char *base, size_t img_size,
 }
 
 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 @@ __tgt_target_table *__tgt_rtl_load_binary_locked(int32_t device_id,
 
     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) {