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

[pdhaliwal]

Summary

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.

[JonChesterfield]

There's one massive object in rtl.cpp that holds a lot of global state. That object would look a lot simpler if it used a vector of struct instead of of multiple vectors, but ultimately there's just a lot of book keeping that currently seems to be necessary. Restricting all globals to a single instance gives us a credible chance of getting init/dtor/locking right.

For the state that can't be eliminated, it probably all needs to end up there. The runtime singleton changed above has been incrementally cut down to a single variable, Environment env_;, and static functions that can be renamed to free functions. I don't remember a reason for leaving env_ in place, likely I had to stop partway through deleting members and do something else.

My preference would be to move (& rename seems good) g_executables out of global scope and into the rtl.cpp object, and delete this runtime class and object entirely. What do you think?

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

47

^ singleton

[pdhaliwal]

Removing Runtime class does make sense. Moving a global out of system.cpp to rtl.cpp will likely trigger more changes, but, that should be manageable.

[pdhaliwal]

Move g_executables to rtl.cpp

[JonChesterfield]

Ok, looks reasonable. There's a comment on the above comment, but could equally just drop the comment. Vector of hsa_executable is probably obvious enough for naming.

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

331

This is per loaded code object, not per device. Kernel information is a bit misleading, it's whatever HSA chose to abstract over the code object.

547

This handling of errors isn't ideal, but matches the current behaviour closely enough

[pdhaliwal]

Removed comment.

[pdhaliwal]

I will submit this to main in my morning.

[JonChesterfield]

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

547

revised this in passing for D102346

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

 /*
  * Machine Info
  */
 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
  */
 
 static hsa_status_t invoke_hsa_copy(hsa_signal_t sig, void *dest,
                                     const void *src, size_t size,
                                     hsa_agent_t agent) {
   const hsa_signal_value_t init = 1;
   const hsa_signal_value_t success = 0;

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

private:
   }
 
   int max_queue_size_;
   int debug_mode_;
 };
 
 class Runtime final {
 public:
   static Runtime &getInstance() {
     static Runtime instance;
     return instance;
   }
 
   // init/finalize
   static atmi_status_t Initialize();
   static atmi_status_t Finalize();
   // machine info
   static atmi_machine_t *GetMachineInfo();
   // modules
   static atmi_status_t RegisterModuleFromMemory(
       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);
   static atmi_status_t Memfree(void *);
   static atmi_status_t Malloc(void **, size_t, atmi_mem_place_t);
 
   int getMaxQueueSize() const { return env_.getMaxQueueSize(); }
   int getDebugMode() const { return env_.getDebugMode(); }

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

 };
 
 // global variables. TODO: Get rid of these
 atmi_machine_t g_atmi_machine;
 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;
 
 bool g_atmi_initialized = false;
 bool g_atmi_hostcall_required = false;
 
 /*
atmi_status_t Runtime::Initialize() {
 
   atl_set_atmi_initialized();
   return ATMI_STATUS_SUCCESS;
 }
 
 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();
   }
   SymbolInfoTable.clear();
   for (uint32_t i = 0; i < KernelInfoTable.size(); i++) {
     KernelInfoTable[i].clear();
   }
   KernelInfoTable.clear();
static hsa_status_t populate_InfoTables(hsa_executable_t executable,
   }
   return HSA_STATUS_SUCCESS;
 }
 
 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);
 
   DEBUG_PRINT("Trying to load module to GPU-%d\n", gpu);
   ATLGPUProcessor &proc = get_processor<ATLGPUProcessor>(place);
   hsa_agent_t agent = proc.agent();
   hsa_executable_t executable = {0};
err = hsa_executable_iterate_symbols(executable, populate_InfoTables,
                                          static_cast<void *>(&gpu));
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Iterating over symbols for execuatable", get_error_string(err));
       exit(1);
     }
 
     // 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;
   }
 }
 
 } // namespace core

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

public:
   int EnvMaxTeamsDefault;
 
   // OpenMP Requires Flags
   int64_t RequiresFlags;
 
   // Resource pools
   SignalPoolT FreeSignalPool;
 
+  std::vector<hsa_executable_t> HSAExecutables;
+
   struct atmiFreePtrDeletor {
     void operator()(void *p) {
       atmi_free(p); // ignore failure to free
     }
   };
 
   // device_State shared across loaded binaries, error if inconsistent size
   std::vector<std::pair<std::unique_ptr<void, atmiFreePtrDeletor>, uint64_t>>
RTLDeviceInfoTy() {
     } else {
       EnvMaxTeamsDefault = -1;
     }
 
     // Default state.
     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
     // atmi_finalize removes access to it
     deviceStateStore.clear();
     KernelArgPoolMap.clear();
     // Terminate hostrpc before finalizing ATMI
     hostrpc_terminate();
+
+    DestroyHSAExecutables();
     atmi_finalize();
   }
 };
 
 pthread_mutex_t SignalPoolT::mutex = PTHREAD_MUTEX_INITIALIZER;
 
 // TODO: May need to drop the trailing to fields until deviceRTL is updated
 struct omptarget_device_environmentTy {
if (rc == 0) {
     *var_size = si.size;
     return ATMI_STATUS_SUCCESS;
   } else {
     return ATMI_STATUS_ERROR;
   }
 }
 
 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
 
 static uint64_t get_device_State_bytes(char *ImageStart, size_t img_size) {
   uint64_t device_State_bytes = 0;
   {
     // If this is the deviceRTL, get the state variable size
     symbol_info size_si;
__tgt_target_table *__tgt_rtl_load_binary_locked(int32_t device_id,
   }
 
   {
     auto env = device_environment(device_id, DeviceInfo.NumberOfDevices, image,
                                   img_size);
 
     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) {
       fprintf(stderr,
               "Possible gpu arch mismatch: device:%s, image:%s please check"
               " compiler flag: -march=<gpu>\n",
               DeviceInfo.GPUName[device_id].c_str(),
               get_elf_mach_gfx_name(elf_e_flags(image)));