[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 ↗

(On Diff #345830)

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.

548 ↗

(On Diff #345830)

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

548 ↗

(On Diff #345830)

revised this in passing for D102346

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

 /*===--------------------------------------------------------------------------
  *              ATMI (Asynchronous Task and Memory Interface)
  *
  * This file is distributed under the MIT License. See LICENSE.txt for details.
  *===------------------------------------------------------------------------*/
 #include "atmi_runtime.h"
 #include "internal.h"
 #include "rt.h"
 #include <hsa.h>
 #include <hsa_ext_amd.h>
 #include <memory>
 
 /*
  * Initialize/Finalize
  */
-atmi_status_t atmi_init() { return core::Runtime::Initialize(); }
+atmi_status_t atmi_init() { return core::Runtime::getInstance().Initialize(); }
 
-atmi_status_t atmi_finalize() { return core::Runtime::Finalize(); }
+atmi_status_t atmi_finalize() {
+  return core::Runtime::getInstance().Finalize();
+}
 
 /*
  * Machine Info
  */
 atmi_machine_t *atmi_machine_get_info() {
   return core::Runtime::GetMachineInfo();
 }
 

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();
+  atmi_status_t Initialize();
+  atmi_status_t Finalize();
   // machine info
   static atmi_machine_t *GetMachineInfo();
   // modules
-  static atmi_status_t RegisterModuleFromMemory(
+  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);
 
   // 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(); }
 
 protected:
   Runtime() = default;
   ~Runtime() = default;
   Runtime(const Runtime &) = delete;
   Runtime &operator=(const Runtime &) = delete;
 
 protected:
   // variable to track environment variables
   Environment env_;
+
+private:
+  std::vector<hsa_executable_t> HSAExecutables;
 };
 
 } // namespace core
 
 #endif // SRC_RUNTIME_INCLUDE_RT_H_

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

 /*===--------------------------------------------------------------------------
  *              ATMI (Asynchronous Task and Memory Interface)
  *
  * This file is distributed under the MIT License. See LICENSE.txt for details.
  *===------------------------------------------------------------------------*/
 #include <gelf.h>
 #include <libelf.h>
 
 #include <cassert>
 #include <cstdarg>
 #include <fstream>
 #include <iomanip>
 #include <iostream>
 #include <set>

 // global variables. TODO: Get rid of these
 atmi_machine_t g_atmi_machine;
 ATLMachine g_atl_machine;
 
 hsa_region_t atl_gpu_kernarg_region;
 std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
 hsa_region_t atl_cpu_kernarg_region;
 
-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]);
+  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));
       exit(1);
     }
   }
 
   for (uint32_t i = 0; i < SymbolInfoTable.size(); i++) {
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