[AMDGPU][Libomptarget] Inline atmi_init/atmi_finalize

[pdhaliwal]

Summary

After D102847, these functions can be inlined.

[JonChesterfield]

Nice, thanks!

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_finalize() { return core::Runtime::Finalize(); }
-
-/*
  * Machine Info
  */
 atmi_machine_t *atmi_machine_get_info() {
   return core::Runtime::GetMachineInfo();
 }
 
 /*
  * Data

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

 #ifndef __cplusplus
 #include <stdbool.h>
 #endif
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
-/** \defgroup context_functions ATMI Context Setup and Finalize
- *  @{
- */
-/**
- * @brief Initialize the ATMI runtime environment.
- *
- * @detal All ATMI runtime functions will fail if this function is not called
- * at least once. The user may initialize difference device types at different
- * regions in the program in order for optimization purposes.
- *
- * @retval ::ATMI_STATUS_SUCCESS The function has executed successfully.
- *
- * @retval ::ATMI_STATUS_ERROR The function encountered errors.
- *
- * @retval ::ATMI_STATUS_UNKNOWN The function encountered errors.
- */
-atmi_status_t atmi_init();
-
-/**
- * @brief Finalize the ATMI runtime environment.
- *
- * @detail ATMI runtime functions will fail if called after finalize.
- *
- * @retval ::ATMI_STATUS_SUCCESS The function has executed successfully.
- *
- * @retval ::ATMI_STATUS_ERROR The function encountered errors.
- *
- * @retval ::ATMI_STATUS_UNKNOWN The function encountered errors.
- */
-atmi_status_t atmi_finalize();
-/** @} */
-
 /** \defgroup module_functions ATMI Module
  * @{
  */
 
 /**
  * @brief Register the ATMI code module from memory on to a specific place
  * (device).
  *

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

 
 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, std::vector<hsa_executable_t> &HSAExecutables);

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

std::vector<ATLGPUProcessor> &gpu_procs =
       g_atl_machine.processors<ATLGPUProcessor>();
   std::vector<hsa_agent_t> agents;
   for (uint32_t i = 0; i < gpu_procs.size(); i++) {
     agents.push_back(gpu_procs[i].agent());
   }
   return hsa_amd_agents_allow_access(agents.size(), &agents[0], NULL, ptr);
 }
 
-atmi_status_t Runtime::Initialize() {
-  atmi_status_t rc = atl_init_gpu_context();
-  if (rc != ATMI_STATUS_SUCCESS) {
-    printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "GPU context init",
-           get_atmi_error_string(atl_init_gpu_context()));
-    return rc;
-  }
-
-  return ATMI_STATUS_SUCCESS;
-}
-
-atmi_status_t Runtime::Finalize() {
-  atmi_status_t rc = ATMI_STATUS_SUCCESS;
-  hsa_status_t err = hsa_shut_down();
-  if (err != HSA_STATUS_SUCCESS) {
-    printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Shutting down HSA",
-           get_error_string(err));
-    rc = ATMI_STATUS_ERROR;
-  }
-
-  return rc;
-}
-
 static void atmi_init_context_structs() {
   atlc.struct_initialized = true; /* This only gets called one time */
   atlc.g_hsa_initialized = false;
   atlc.g_gpu_initialized = false;
   atlc.g_tasks_initialized = false;
 }
 
 // Implement memory_pool iteration function

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

RTLDeviceInfoTy() {
     //  1 => tracing dispatch only
     // >1 => verbosity increase
     if (char *envStr = getenv("LIBOMPTARGET_KERNEL_TRACE"))
       print_kernel_trace = atoi(envStr);
     else
       print_kernel_trace = 0;
 
     DP("Start initializing HSA-ATMI\n");
-    atmi_status_t err = atmi_init();
+    atmi_status_t err = core::atl_init_gpu_context();
     if (err != ATMI_STATUS_SUCCESS) {
       DP("Error when initializing HSA-ATMI\n");
       return;
     }
+
     // Init hostcall soon after initializing ATMI
     hostrpc_init();
 
     HSAAgents = find_gpu_agents();
     NumberOfDevices = (int)HSAAgents.size();
 
     if (NumberOfDevices == 0) {
       DP("There are no devices supporting HSA.\n");
~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();
 
+    hsa_status_t Err;
     for (uint32_t I = 0; I < HSAExecutables.size(); I++) {
-      hsa_status_t Err = hsa_executable_destroy(HSAExecutables[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));
       }
     }
 
-    atmi_finalize();
+    Err = hsa_shut_down();
+    if (Err != HSA_STATUS_SUCCESS) {
+      printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Shutting down HSA",
+             get_error_string(Err));
+    }
   }
 };
 
 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 {
   int32_t debug_level; // gets value of envvar LIBOMPTARGET_DEVICE_RTL_DEBUG