[libomptarget][amdgpu] Precisely manage hsa lifetime

[JonChesterfield]

Summary

The hsa library must be initialized before any calls into it and
destructed after the last call into it. There have been a number of bugs in
this area related to member variables which would like to use raii to manage
resources acquired from hsa.

This patch moves the init/shutdown of hsa into a class, such that when used as
the first member variable (could be a base), the lifetime of other member
variables are reliably scoped within it. This will allow other classes to use
raii reliably when used as member variables within the global.

[pdhaliwal]

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

846

This can be removed.

[JonChesterfield]

• Drop second shutdown call

[JonChesterfield]

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

846

Thanks!

[pdhaliwal]

Looks good! Thanks.

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

 //===----------------------------------------------------------------------===//
 
 #include <algorithm>
 #include <assert.h>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <functional>
 #include <libelf.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <shared_mutex>
 #include <unordered_map>
 #include <vector>
 
 #include "interop_hsa.h"
struct EnvironmentVariables {
   int MaxTeamsDefault;
 };
 
 template <uint32_t wavesize>
 static constexpr const llvm::omp::GV &getGridValue() {
   return llvm::omp::getAMDGPUGridValues<wavesize>();
 }
 
+struct HSALifetime {
+  // Wrapper around HSA used to ensure it is constructed before other types
+  // and destructed after, which means said other types can use raii for
+  // cleanup without risking running outside of the lifetime of HSA
+  const hsa_status_t S;
+
+  bool success() { return S == HSA_STATUS_SUCCESS; }
+  HSALifetime() : S(hsa_init()) {}
+
+  ~HSALifetime() {
+    if (S == HSA_STATUS_SUCCESS) {
+      hsa_status_t Err = hsa_shut_down();
+      if (Err != HSA_STATUS_SUCCESS) {
+        // Can't call into HSA to get a string from the integer
+        DP("Shutting down HSA failed: %d\n", Err);
+      }
+    }
+  }
+};
+
 /// Class containing all the device information
 class RTLDeviceInfoTy {
+  HSALifetime HSA; // First field => constructed first and destructed last
   std::vector<std::list<FuncOrGblEntryTy>> FuncGblEntries;
-  bool HSAInitializeSucceeded = false;
 
 public:
   // load binary populates symbol tables and mutates various global state
   // run uses those symbol tables
   std::shared_timed_mutex load_run_lock;
 
   int NumberOfDevices = 0;
 
static int readEnvElseMinusOne(const char *Env) {
     if (envStr) {
       res = std::stoi(envStr);
       DP("Parsed %s=%d\n", Env, res);
     }
     return res;
   }
 
   RTLDeviceInfoTy() {
+    DP("Start initializing " GETNAME(TARGET_NAME) "\n");
+
     // LIBOMPTARGET_KERNEL_TRACE provides a kernel launch trace to stderr
     // anytime. You do not need a debug library build.
     //  0 => no tracing
     //  1 => tracing dispatch only
     // >1 => verbosity increase
+
+    if (!HSA.success()) {
+      DP("Error when initializing HSA in " GETNAME(TARGET_NAME) "\n");
+      return;
+    }
+
     if (char *envStr = getenv("LIBOMPTARGET_KERNEL_TRACE"))
       print_kernel_trace = atoi(envStr);
     else
       print_kernel_trace = 0;
 
-    DP("Start initializing " GETNAME(TARGET_NAME) "\n");
-    hsa_status_t err = hsa_init();
-    if (err == HSA_STATUS_SUCCESS) {
-      err = core::atl_init_gpu_context();
-    }
-    if (err == HSA_STATUS_SUCCESS) {
-      HSAInitializeSucceeded = true;
-    } else {
+    hsa_status_t err = core::atl_init_gpu_context();
+    if (err != HSA_STATUS_SUCCESS) {
       DP("Error when initializing " GETNAME(TARGET_NAME) "\n");
       return;
     }
 
     // Init hostcall soon after initializing hsa
     hostrpc_init();
 
     err = FindAgents([&](hsa_device_type_t DeviceType, hsa_agent_t Agent) {
RTLDeviceInfoTy() {
     Env.TeamThreadLimit = readEnvElseMinusOne("OMP_TEAMS_THREAD_LIMIT");
 
     // Default state.
     RequiresFlags = OMP_REQ_UNDEFINED;
   }
 
   ~RTLDeviceInfoTy() {
     DP("Finalizing the " GETNAME(TARGET_NAME) " DeviceInfo.\n");
-    if (!HSAInitializeSucceeded) {
+    if (!HSA.success()) {
       // Then none of these can have been set up and they can't be torn down
       return;
     }
     // Run destructors on types that use HSA before
     // impl_finalize removes access to it
     deviceStateStore.clear();
     KernelArgPoolMap.clear();
     // Terminate hostrpc before finalizing hsa
     hostrpc_terminate();
 
     hsa_status_t Err;
     for (uint32_t I = 0; I < HSAExecutables.size(); I++) {
       Err = hsa_executable_destroy(HSAExecutables[I]);
       if (Err != HSA_STATUS_SUCCESS) {
         DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
            "Destroying executable", get_error_string(Err));
       }
     }
 
     Err = hsa_shut_down();
     if (Err != HSA_STATUS_SUCCESS) {
       DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Shutting down HSA",
          get_error_string(Err));
     }
   }
 };
 
 pthread_mutex_t SignalPoolT::mutex = PTHREAD_MUTEX_INITIALIZER;