diff --git a/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp b/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp --- a/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp +++ b/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp @@ -10,13 +10,6 @@ #include #include -/* - * 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 --- 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 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 * @{ */ diff --git a/openmp/libomptarget/plugins/amdgpu/impl/rt.h b/openmp/libomptarget/plugins/amdgpu/impl/rt.h --- a/openmp/libomptarget/plugins/amdgpu/impl/rt.h +++ b/openmp/libomptarget/plugins/amdgpu/impl/rt.h @@ -49,9 +49,6 @@ 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 --- a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp +++ b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp @@ -171,29 +171,6 @@ 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 --- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp +++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp @@ -466,11 +466,12 @@ 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 @@ // 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)); + } } };