[Openmp-commits] [openmp] d7503c3 - [AMDGPU][Libomptarget] Rename & move g_executables to private
Pushpinder Singh via Openmp-commits
openmp-commits at lists.llvm.org
Mon May 17 22:43:33 PDT 2021
Author: Pushpinder Singh
Date: 2021-05-18T05:43:23Z
New Revision: d7503c3bce491e2672386906ec879c7df5ede7a5
URL: https://github.com/llvm/llvm-project/commit/d7503c3bce491e2672386906ec879c7df5ede7a5
DIFF: https://github.com/llvm/llvm-project/commit/d7503c3bce491e2672386906ec879c7df5ede7a5.diff
LOG: [AMDGPU][Libomptarget] Rename & move g_executables to private
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.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D102600
Added:
Modified:
openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
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 285dc2dbe763..76e0bc6da194 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/atmi.cpp
@@ -24,18 +24,6 @@ atmi_machine_t *atmi_machine_get_info() {
return core::Runtime::GetMachineInfo();
}
-/*
- * Modules
- */
-atmi_status_t atmi_module_register_from_memory_to_place(
- void *module_bytes, size_t module_size, atmi_place_t place,
- atmi_status_t (*on_deserialized_data)(void *data, size_t size,
- void *cb_state),
- void *cb_state) {
- return core::Runtime::getInstance().RegisterModuleFromMemory(
- module_bytes, module_size, place, on_deserialized_data, cb_state);
-}
-
/*
* Data
*/
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/rt.h b/openmp/libomptarget/plugins/amdgpu/impl/rt.h
index a857861307c6..9954725c67a5 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/rt.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/rt.h
@@ -59,7 +59,7 @@ class Runtime final {
void *, size_t, atmi_place_t,
atmi_status_t (*on_deserialized_data)(void *data, size_t size,
void *cb_state),
- void *cb_state);
+ void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
// data
static atmi_status_t Memcpy(hsa_signal_t, void *, const void *, size_t);
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
index 5e45bdba56a7..02d04e913346 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
@@ -132,8 +132,6 @@ ATLMachine g_atl_machine;
std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
-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;
@@ -204,15 +202,6 @@ atmi_status_t Runtime::Initialize() {
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]);
- 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++) {
SymbolInfoTable[i].clear();
}
@@ -1170,7 +1159,7 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
void *module_bytes, size_t module_size, atmi_place_t place,
atmi_status_t (*on_deserialized_data)(void *data, size_t size,
void *cb_state),
- void *cb_state) {
+ void *cb_state, std::vector<hsa_executable_t> &HSAExecutables) {
hsa_status_t err;
int gpu = place.device_id;
assert(gpu >= 0);
@@ -1269,7 +1258,7 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
}
// 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;
diff --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
index 8c50ac388d9c..944d9d6dd59c 100644
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -328,6 +328,8 @@ class RTLDeviceInfoTy {
// Resource pools
SignalPoolT FreeSignalPool;
+ std::vector<hsa_executable_t> HSAExecutables;
+
struct atmiFreePtrDeletor {
void operator()(void *p) {
atmi_free(p); // ignore failure to free
@@ -538,6 +540,18 @@ class RTLDeviceInfoTy {
RequiresFlags = OMP_REQ_UNDEFINED;
}
+ void DestroyHSAExecutables() {
+ hsa_status_t Err;
+ 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));
+ return;
+ }
+ }
+ }
+
~RTLDeviceInfoTy() {
DP("Finalizing the HSA-ATMI DeviceInfo.\n");
// Run destructors on types that use HSA before
@@ -546,6 +560,8 @@ class RTLDeviceInfoTy {
KernelArgPoolMap.clear();
// Terminate hostrpc before finalizing ATMI
hostrpc_terminate();
+
+ DestroyHSAExecutables();
atmi_finalize();
}
};
@@ -971,15 +987,16 @@ atmi_status_t interop_get_symbol_info(char *base, size_t img_size,
}
template <typename C>
-atmi_status_t module_register_from_memory_to_place(void *module_bytes,
- size_t module_size,
- atmi_place_t place, C cb) {
+atmi_status_t module_register_from_memory_to_place(
+ void *module_bytes, size_t module_size, atmi_place_t place, C cb,
+ std::vector<hsa_executable_t> &HSAExecutables) {
auto L = [](void *data, size_t size, void *cb_state) -> atmi_status_t {
C *unwrapped = static_cast<C *>(cb_state);
return (*unwrapped)(data, size);
};
- return atmi_module_register_from_memory_to_place(
- module_bytes, module_size, place, L, static_cast<void *>(&cb));
+ return core::Runtime::RegisterModuleFromMemory(
+ module_bytes, module_size, place, L, static_cast<void *>(&cb),
+ HSAExecutables);
}
} // namespace
@@ -1180,9 +1197,8 @@ __tgt_target_table *__tgt_rtl_load_binary_locked(int32_t device_id,
atmi_status_t err = module_register_from_memory_to_place(
(void *)image->ImageStart, img_size, get_gpu_place(device_id),
- [&](void *data, size_t size) {
- return env.before_loading(data, size);
- });
+ [&](void *data, size_t size) { return env.before_loading(data, size); },
+ DeviceInfo.HSAExecutables);
check("Module registering", err);
if (err != ATMI_STATUS_SUCCESS) {
More information about the Openmp-commits
mailing list