[Openmp-commits] [openmp] a2d6ef5 - [AMDGPU][Libomptarget] Inline atmi_init/atmi_finalize

[Pushpinder Singh via Openmp-commits]

Author: Pushpinder Singh
Date: 2021-05-26T10:50:08Z
New Revision: a2d6ef58765301fa95776cd17033a0974a487bf4

URL: https://github.com/llvm/llvm-project/commit/a2d6ef58765301fa95776cd17033a0974a487bf4
DIFF: https://github.com/llvm/llvm-project/commit/a2d6ef58765301fa95776cd17033a0974a487bf4.diff

LOG: [AMDGPU][Libomptarget] Inline atmi_init/atmi_finalize

After D102847, these functions can be inlined.

Reviewed By: JonChesterfield

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

Added: 
    

Modified: 
    openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
    openmp/libomptarget/plugins/amdgpu/impl/atmi_runtime.h
    openmp/libomptarget/plugins/amdgpu/impl/rt.h
    openmp/libomptarget/plugins/amdgpu/impl/system.cpp
    openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp b/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
index 76e0bc6da194..985f1ba6a5b0 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
@@ -10,13 +10,6 @@
 #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
  */

diff  --git a/openmp/libomptarget/plugins/amdgpu/impl/atmi_runtime.h b/openmp/libomptarget/plugins/amdgpu/impl/atmi_runtime.h
index 47022f7f5dea..20919a4cd9f3 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/atmi_runtime.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/atmi_runtime.h
@@ -18,38 +18,6 @@
 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 
diff erence device types at 
diff erent
- * 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
  * @{
  */

diff  --git a/openmp/libomptarget/plugins/amdgpu/impl/rt.h b/openmp/libomptarget/plugins/amdgpu/impl/rt.h
index 9954725c67a5..622f2d434817 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/rt.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/rt.h
@@ -49,9 +49,6 @@ class Runtime final {
     return instance;
   }
 
-  // init/finalize
-  static atmi_status_t Initialize();
-  static atmi_status_t Finalize();
   // machine info
   static atmi_machine_t *GetMachineInfo();
   // modules

diff  --git a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
index 08168da9e838..a2d9e3bcf962 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
@@ -171,29 +171,6 @@ hsa_status_t allow_access_to_all_gpu_agents(void *ptr) {
   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;

diff  --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
index 4883288e0725..9ac1d806ef9f 100644
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -466,11 +466,12 @@ class RTLDeviceInfoTy {
       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();
 
@@ -574,15 +575,20 @@ class RTLDeviceInfoTy {
     // 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));
+    }
   }
 };