From 304568c1a314299050c53fbcd80a123392c125e7 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Wed, 5 Sep 2018 15:49:33 -0500 Subject: [PATCH] kernel free on executable_destroy Change-Id: I062e08fd5df94f4d21b443b74494dc6e95201f3d --- inc/rocprofiler.h | 1 + src/core/hsa_interceptor.h | 64 +++++++++++++++++++++++++++----------- src/core/rocprofiler.cpp | 1 + test/tool/tool.cpp | 1 + 4 files changed, 49 insertions(+), 18 deletions(-) diff --git a/inc/rocprofiler.h b/inc/rocprofiler.h index 8e198231ec..7b52c2aefb 100644 --- a/inc/rocprofiler.h +++ b/inc/rocprofiler.h @@ -479,6 +479,7 @@ struct rocprofiler_hsa_callback_data_t { size_t size; // allocated area size, zero size means 'free' callback hsa_amd_segment_t segment; // allocated area's memory segment type hsa_amd_memory_pool_global_flag_t global_flag; // allocated area's memory global flag + int is_code; // equal to 1 if code is allocated } allocate; struct { hsa_device_type_t type; // type of assigned device diff --git a/src/core/hsa_interceptor.h b/src/core/hsa_interceptor.h index 7027018571..0f285d49f4 100644 --- a/src/core/hsa_interceptor.h +++ b/src/core/hsa_interceptor.h @@ -67,6 +67,7 @@ extern decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn; extern decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn; extern decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn; extern decltype(hsa_executable_freeze)* hsa_executable_freeze_fn; +extern decltype(hsa_executable_destroy)* hsa_executable_destroy_fn; class HsaInterceptor { public: @@ -93,6 +94,7 @@ class HsaInterceptor { hsa_amd_agents_allow_access_fn = table->amd_ext_->hsa_amd_agents_allow_access_fn; hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn; hsa_executable_freeze_fn = table->core_->hsa_executable_freeze_fn; + hsa_executable_destroy_fn = table->core_->hsa_executable_destroy_fn; // Intercepting HSA API table->core_->hsa_memory_allocate_fn = MemoryAllocate; @@ -103,6 +105,7 @@ class HsaInterceptor { table->amd_ext_->hsa_amd_agents_allow_access_fn = AgentsAllowAccess; table->amd_ext_->hsa_amd_memory_async_copy_fn = MemoryAsyncCopy; table->core_->hsa_executable_freeze_fn = ExecutableFreeze; + table->core_->hsa_executable_destroy_fn = ExecutableDestroy; } } @@ -284,8 +287,9 @@ class HsaInterceptor { static hsa_status_t CodeObjectCallback( hsa_executable_t executable, hsa_loaded_code_object_t loaded_code_object, - void *) + void* arg) { + const int free_flag = reinterpret_cast(arg); rocprofiler_hsa_callback_data_t data{}; HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info( @@ -293,31 +297,38 @@ class HsaInterceptor { HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE, &data.allocate.ptr)); - HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, - HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE, - &data.allocate.size)); + if (free_flag == 0) { + HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info( + loaded_code_object, + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE, + &data.allocate.size)); + } else { + data.allocate.size = 0; + } // Local GPU memory // GLOBAL; FLAGS: COARSE GRAINED data.allocate.segment = HSA_AMD_SEGMENT_GLOBAL; data.allocate.global_flag = HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED; + data.allocate.is_code = 1; ISSUE_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE); - IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) { - hsa_amd_pointer_info_t pointer_info{}; - uint32_t num_agents = 0; - hsa_agent_t* agents = NULL; - pointer_info.size = sizeof(hsa_amd_pointer_info_t); - HSA_RT(hsa_amd_pointer_info( - const_cast(data.allocate.ptr), - &pointer_info, - malloc, - &num_agents, - &agents)); + if (free_flag == 0) { + IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) { + hsa_amd_pointer_info_t pointer_info{}; + uint32_t num_agents = 0; + hsa_agent_t* agents = NULL; + pointer_info.size = sizeof(hsa_amd_pointer_info_t); + HSA_RT(hsa_amd_pointer_info( + const_cast(data.allocate.ptr), + &pointer_info, + malloc, + &num_agents, + &agents)); - DeviceCallback(num_agents, agents, data.allocate.ptr); + DeviceCallback(num_agents, agents, data.allocate.ptr); + } } return HSA_STATUS_SUCCESS; @@ -335,12 +346,29 @@ class HsaInterceptor { LoaderApiTable.hsa_ven_amd_loader_executable_iterate_loaded_code_objects( executable, CodeObjectCallback, - NULL); + reinterpret_cast(0)); } return status; } + static hsa_status_t ExecutableDestroy( + hsa_executable_t executable) + { + hsa_status_t status = HSA_STATUS_SUCCESS; + + IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) { + LoaderApiTable.hsa_ven_amd_loader_executable_iterate_loaded_code_objects( + executable, + CodeObjectCallback, + reinterpret_cast(1)); + } + + HSA_RT(hsa_executable_destroy_fn(executable)); + + return status; + } + static bool enable_; static hsa_ven_amd_loader_1_01_pfn_t LoaderApiTable; static rocprofiler_hsa_callbacks_t callbacks_; diff --git a/src/core/rocprofiler.cpp b/src/core/rocprofiler.cpp index 70c6ea0c3b..2fe9777b0e 100644 --- a/src/core/rocprofiler.cpp +++ b/src/core/rocprofiler.cpp @@ -94,6 +94,7 @@ decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn; decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn; decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn; decltype(hsa_executable_freeze)* hsa_executable_freeze_fn; +decltype(hsa_executable_destroy)* hsa_executable_destroy_fn; ::HsaApiTable* kHsaApiTable; diff --git a/test/tool/tool.cpp b/test/tool/tool.cpp index 34bd0516e9..1dfbe5a9a0 100644 --- a/test/tool/tool.cpp +++ b/test/tool/tool.cpp @@ -842,6 +842,7 @@ hsa_status_t hsa_unified_callback( printf(" alloc size = %zu\n", data->allocate.size); printf(" segment type = 0x%x\n", data->allocate.segment); printf(" global flag = 0x%x\n", data->allocate.global_flag); + printf(" is_code = %x\n", data->allocate.is_code); break; case ROCPROFILER_HSA_CB_ID_DEVICE: printf(" device type = 0x%x\n", data->device.type);