diff --git a/script/hsaap.py b/script/hsaap.py index 7139be879f..dab2137c2f 100755 --- a/script/hsaap.py +++ b/script/hsaap.py @@ -29,10 +29,10 @@ H_OUT='hsa_prof_str.h' CPP_OUT='hsa_prof_str.inline.h' API_TABLES_H = 'hsa_api_trace.h' API_HEADERS_H = ( - ('CoreApiTable', 'hsa.h'), - ('AmdExtTable', 'hsa_ext_amd.h'), - ('ImageExtTable', 'hsa_ext_image.h'), - ('AmdExtTable', API_TABLES_H), + ('CoreApi', 'hsa.h'), + ('AmdExt', 'hsa_ext_amd.h'), + ('ImageExt', 'hsa_ext_image.h'), + ('AmdExt', API_TABLES_H), ) LICENSE = \ @@ -106,7 +106,7 @@ class API_TableParser: self.inp = open(header, 'r') - self.beg_pattern = re.compile('^\s*struct\s+' + name + '\s*{\s*$') + self.beg_pattern = re.compile('^\s*struct\s+' + name + 'Table\s*{\s*$') self.end_pattern = re.compile('^\s*};\s*$') self.array = [] self.parse() @@ -330,9 +330,9 @@ class API_DescrParser: self.cpp_content += 'namespace roctracer {\n' self.cpp_content += 'namespace hsa_support {\n\n' - self.cpp_content += 'static CoreApiTable CoreApiTable_saved;\n' - self.cpp_content += 'static AmdExtTable AmdExtTable_saved;\n' - self.cpp_content += 'static ImageExtTable ImageExtTable_saved;\n\n' + self.cpp_content += 'static CoreApiTable CoreApi_saved_before_cb;\n' + self.cpp_content += 'static AmdExtTable AmdExt_saved_before_cb;\n' + self.cpp_content += 'static ImageExtTable ImageExt_saved_before_cb;\n\n' self.cpp_content += 'std::atomic hsa_counter_{1};\n' self.cpp_content += 'static thread_local uint64_t hsa_correlation_id_tls = 0;\n' @@ -428,7 +428,7 @@ class API_DescrParser: content += ' if (api_callback_fun) api_callback_fun(ACTIVITY_DOMAIN_HSA_API, ' + call_id + ', &api_data, api_callback_arg);\n' if ret_type != 'void': content += ' ' + ret_type + ' ret =' - content += ' ' + name + '_saved.' + call + '_fn(' + ', '.join(struct['alst']) + ');\n' + content += ' ' + name + '_saved_before_cb.' + call + '_fn(' + ', '.join(struct['alst']) + ');\n' if ret_type != 'void': content += ' api_data.' + ret_type + '_retval = ret;\n' content += ' api_data.phase = 1;\n' @@ -444,8 +444,8 @@ class API_DescrParser: if n > 0 and call == '-': content += '};\n' if n == 0 or (call == '-' and name != '-'): - content += 'static void intercept_' + name + '(' + name + '* table) {\n' - content += ' ' + name + '_saved = *table;\n' + content += 'static void Install' + name + 'Wrappers(' + name + 'Table* table) {\n' + content += ' ' + name + '_saved_before_cb = *table;\n' if call != '-': if call != 'hsa_shut_down': content += ' table->' + call + '_fn = ' + call + '_callback;\n' diff --git a/src/roctracer/roctracer.cpp b/src/roctracer/roctracer.cpp index c7630def9b..b4415ae0e4 100644 --- a/src/roctracer/roctracer.cpp +++ b/src/roctracer/roctracer.cpp @@ -106,28 +106,12 @@ mark_api_callback_t* mark_api_callback_ptr = nullptr; namespace roctracer { namespace hsa_support { -decltype(hsa_system_get_info)* hsa_system_get_info_fn = hsa_system_get_info; -decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn = hsa_amd_memory_async_copy; -decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn = - hsa_amd_memory_async_copy_rect; -::HsaApiTable* kHsaApiTable; - -void SaveHsaApi(::HsaApiTable* table) { - kHsaApiTable = table; - hsa_system_get_info_fn = table->core_->hsa_system_get_info_fn; - hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn; - hsa_amd_memory_async_copy_rect_fn = table->amd_ext_->hsa_amd_memory_async_copy_rect_fn; -} - -void RestoreHsaApi() { - ::HsaApiTable* table = kHsaApiTable; - table->core_->hsa_system_get_info_fn = hsa_system_get_info_fn; - table->amd_ext_->hsa_amd_memory_async_copy_fn = hsa_amd_memory_async_copy_fn; - table->amd_ext_->hsa_amd_memory_async_copy_rect_fn = hsa_amd_memory_async_copy_rect_fn; -} +static CoreApiTable saved_core_api; +static AmdExtTable saved_amd_ext_api; // async copy activity callback +std::mutex init_mutex; bool async_copy_callback_enabled = false; MemoryPool* async_copy_callback_memory_pool = nullptr; @@ -143,14 +127,15 @@ namespace util { roctracer_timestamp_t timestamp_ns() { uint64_t sysclock; - hsa_status_t status = hsa_support::hsa_system_get_info_fn(HSA_SYSTEM_INFO_TIMESTAMP, &sysclock); + hsa_status_t status = + hsa_support::saved_core_api.hsa_system_get_info_fn(HSA_SYSTEM_INFO_TIMESTAMP, &sysclock); if (status == HSA_STATUS_ERROR_NOT_INITIALIZED) return 0; CHECK_HSA_STATUS("hsa_system_get_info()", status); static uint64_t sysclock_period = []() { uint64_t sysclock_hz = 0; - hsa_status_t status = - hsa_support::hsa_system_get_info_fn(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz); + hsa_status_t status = hsa_support::saved_core_api.hsa_system_get_info_fn( + HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz); CHECK_HSA_STATUS("hsa_system_get_info()", status); return (uint64_t)1000000000 / sysclock_hz; }(); @@ -447,20 +432,21 @@ hsa_status_t hsa_amd_memory_async_copy_interceptor(void* dst, hsa_agent_t dst_ag size_t size, uint32_t num_dep_signals, const hsa_signal_t* dep_signals, hsa_signal_t completion_signal) { - hsa_status_t status = HSA_STATUS_SUCCESS; - if (hsa_support::async_copy_callback_enabled) { - Tracker::entry_t* entry = new Tracker::entry_t(); - entry->handler = hsa_async_copy_handler; - entry->pool = hsa_support::async_copy_callback_memory_pool; - entry->correlation_id = hsa_correlation_id_tls; - Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry); - status = hsa_amd_memory_async_copy_fn(dst, dst_agent, src, src_agent, size, num_dep_signals, - dep_signals, entry->signal); - if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry); - } else { - status = hsa_amd_memory_async_copy_fn(dst, dst_agent, src, src_agent, size, num_dep_signals, - dep_signals, completion_signal); + if (!async_copy_callback_enabled) { + return saved_amd_ext_api.hsa_amd_memory_async_copy_fn( + dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, completion_signal); } + + Tracker::entry_t* entry = new Tracker::entry_t(); + entry->handler = hsa_async_copy_handler; + entry->pool = async_copy_callback_memory_pool; + entry->correlation_id = hsa_correlation_id_tls; + Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry); + + hsa_status_t status = saved_amd_ext_api.hsa_amd_memory_async_copy_fn( + dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, entry->signal); + if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry); + return status; } @@ -469,21 +455,23 @@ hsa_status_t hsa_amd_memory_async_copy_rect_interceptor( const hsa_dim3_t* src_offset, const hsa_dim3_t* range, hsa_agent_t copy_agent, hsa_amd_copy_direction_t dir, uint32_t num_dep_signals, const hsa_signal_t* dep_signals, hsa_signal_t completion_signal) { - hsa_status_t status = HSA_STATUS_SUCCESS; - if (hsa_support::async_copy_callback_enabled) { - Tracker::entry_t* entry = new Tracker::entry_t(); - entry->handler = hsa_async_copy_handler; - entry->pool = hsa_support::async_copy_callback_memory_pool; - entry->correlation_id = hsa_correlation_id_tls; - Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry); - status = hsa_amd_memory_async_copy_rect_fn(dst, dst_offset, src, src_offset, range, copy_agent, - dir, num_dep_signals, dep_signals, entry->signal); - if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry); - } else { - status = - hsa_amd_memory_async_copy_rect_fn(dst, dst_offset, src, src_offset, range, copy_agent, dir, - num_dep_signals, dep_signals, completion_signal); + if (!async_copy_callback_enabled) { + return saved_amd_ext_api.hsa_amd_memory_async_copy_rect_fn( + dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals, + completion_signal); } + + Tracker::entry_t* entry = new Tracker::entry_t(); + entry->handler = hsa_async_copy_handler; + entry->pool = async_copy_callback_memory_pool; + entry->correlation_id = hsa_correlation_id_tls; + Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry); + + hsa_status_t status = saved_amd_ext_api.hsa_amd_memory_async_copy_rect_fn( + dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals, + entry->signal); + if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry); + return status; } @@ -855,6 +843,13 @@ static void roctracer_enable_activity_fun(roctracer_domain_t domain, uint32_t op switch (domain) { case ACTIVITY_DOMAIN_HSA_OPS: { if (op == HSA_OP_ID_COPY) { + std::scoped_lock lock(hsa_support::init_mutex); + + if (hsa_support::saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn != nullptr) { + [[maybe_unused]] hsa_status_t status = + hsa_support::saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(true); + assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed"); + } RocpLoader::Instance(); hsa_support::async_copy_callback_enabled = true; hsa_support::async_copy_callback_memory_pool = reinterpret_cast(pool); @@ -976,8 +971,17 @@ static void roctracer_disable_activity_fun(roctracer_domain_t domain, uint32_t o switch (domain) { case ACTIVITY_DOMAIN_HSA_OPS: { if (op == HSA_OP_ID_COPY) { + std::scoped_lock lock(hsa_support::init_mutex); + hsa_support::async_copy_callback_enabled = false; hsa_support::async_copy_callback_memory_pool = nullptr; + + if (hsa_support::saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn != nullptr) { + [[maybe_unused]] hsa_status_t status = + hsa_support::saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(false); + assert(status == HSA_STATUS_SUCCESS || status == HSA_STATUS_ERROR_NOT_INITIALIZED || + !"hsa_amd_profiling_async_copy_enable failed"); + } } else { if (RocpLoader::GetRef() == nullptr) break; if (!RocpLoader::Instance().EnableActivityCallback(op, false)) @@ -1178,33 +1182,12 @@ ROCTRACER_API roctracer_status_t roctracer_set_properties(roctracer_domain_t dom API_METHOD_PREFIX switch (domain) { case ACTIVITY_DOMAIN_HSA_OPS: { - // HSA OPS properties - hsa_ops_properties_t* ops_properties = reinterpret_cast(properties); - HsaApiTable* table = reinterpret_cast(ops_properties->table); - - // HSA async-copy tracing - [[maybe_unused]] hsa_status_t status = hsa_amd_profiling_async_copy_enable(true); - assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed"); - hsa_support::hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn; - hsa_support::hsa_amd_memory_async_copy_rect_fn = - table->amd_ext_->hsa_amd_memory_async_copy_rect_fn; - table->amd_ext_->hsa_amd_memory_async_copy_fn = - hsa_support::hsa_amd_memory_async_copy_interceptor; - table->amd_ext_->hsa_amd_memory_async_copy_rect_fn = - hsa_support::hsa_amd_memory_async_copy_rect_interceptor; - - break; } case ACTIVITY_DOMAIN_HSA_EVT: { break; } case ACTIVITY_DOMAIN_HSA_API: { - // HSA API properties - HsaApiTable* table = reinterpret_cast(properties); - hsa_support::intercept_CoreApiTable(table->core_); - hsa_support::intercept_AmdExtTable(table->amd_ext_); - hsa_support::intercept_ImageExtTable(table->image_ext_); break; } case ACTIVITY_DOMAIN_HIP_OPS: @@ -1246,10 +1229,38 @@ ROCTRACER_EXPORT extern const uint32_t HSA_AMD_TOOL_PRIORITY = 50; // HSA-runtime tool on-load method ROCTRACER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count, const char* const* failed_tool_names) { - hsa_support::SaveHsaApi(table); + std::scoped_lock lock(hsa_support::init_mutex); + + // Save the HSA core api and amd_ext api. + hsa_support::saved_core_api = *table->core_; + hsa_support::saved_amd_ext_api = *table->amd_ext_; + + // Install the HSA_OPS intercept + table->amd_ext_->hsa_amd_memory_async_copy_fn = + hsa_support::hsa_amd_memory_async_copy_interceptor; + table->amd_ext_->hsa_amd_memory_async_copy_rect_fn = + hsa_support::hsa_amd_memory_async_copy_rect_interceptor; + + // Install the HSA_API wrappers + hsa_support::InstallCoreApiWrappers(table->core_); + hsa_support::InstallAmdExtWrappers(table->amd_ext_); + hsa_support::InstallImageExtWrappers(table->image_ext_); + + if (hsa_support::async_copy_callback_enabled) { + [[maybe_unused]] hsa_status_t status = + hsa_support::saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(true); + assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed"); + } + return true; } -ROCTRACER_EXPORT void OnUnload() {} +ROCTRACER_EXPORT void OnUnload() { + if (hsa_support::async_copy_callback_enabled) { + [[maybe_unused]] hsa_status_t status = + hsa_support::saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(false); + assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed"); + } +} } // extern "C" \ No newline at end of file diff --git a/src/tracer_tool/tracer_tool.cpp b/src/tracer_tool/tracer_tool.cpp index ce5f555f80..f627bf5f41 100644 --- a/src/tracer_tool/tracer_tool.cpp +++ b/src/tracer_tool/tracer_tool.cpp @@ -931,9 +931,6 @@ ROCTRACER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, if (trace_hsa_api) { hsa_api_file_handle = open_output_file(output_prefix, "hsa_api_trace.txt"); - // initialize HSA tracing - roctracer_set_properties(ACTIVITY_DOMAIN_HSA_API, (void*)table); - fprintf(stdout, " HSA-trace("); fflush(stdout); if (hsa_api_vec.size() != 0) { @@ -956,11 +953,6 @@ ROCTRACER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, if (trace_hsa_activity) { hsa_async_copy_file_handle = open_output_file(output_prefix, "async_copy_trace.txt"); - // initialize HSA tracing - roctracer::hsa_ops_properties_t ops_properties{}; - ops_properties.table = table; - roctracer_set_properties(ACTIVITY_DOMAIN_HSA_OPS, &ops_properties); - // Allocating tracing pool open_tracing_pool(); diff --git a/test/golden_traces/copy_hsa_trace.txt b/test/golden_traces/copy_hsa_trace.txt index a97c7c354b..be1e2dd18f 100644 --- a/test/golden_traces/copy_hsa_trace.txt +++ b/test/golden_traces/copy_hsa_trace.txt @@ -1,69 +1,68 @@ -ROCTracer (pid=882613): -0x555e40292a90 agent cpu -0x555e40295470 agent gpu -0x555e402cd500 agent gpu -538166319764062 +ROCTracer (pid=566828): +0x55e1b9d507c0 agent cpu +0x55e1b9d4eeb0 agent gpu +0x55e1b9d8b540 agent gpu +975779239309496 HSA-trace() HSA-activity-trace() -538166319915287:538166319920086 882613:882613 hsa_amd_profiling_async_copy_enable(1) = 0 -538166320173874:538166320174335 882613:882613 hsa_agent_get_info({handle=93863291726480}, 0, 0x7ffddcd49dd8) = 0 -538166320175287:538166320175548 882613:882613 hsa_agent_get_info({handle=93863291726480}, 17, 0x7ffddcd49dc0) = 0 -538166320177511:538166320178523 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291726976}, 0, 0x7ffddcd49bdc) = 0 -538166320179144:538166320179415 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291726976}, 1, 0x7ffddcd49be0) = 0 -538166320180337:538166320180587 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291726976}, 2, 0x7ffddcd49c00) = 0 -538166320181198:538166320181449 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291726976}, 6, 0x7ffddcd49c08) = 0 -538166320183613:538166320183873 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291727280}, 0, 0x7ffddcd49bdc) = 0 -538166320184484:538166320184745 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291727280}, 1, 0x7ffddcd49be0) = 0 -538166320185336:538166320185587 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291727280}, 2, 0x7ffddcd49c00) = 0 -538166320186178:538166320186428 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291727280}, 6, 0x7ffddcd49c08) = 0 -538166320187390:538166320188031 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291728112}, 0, 0x7ffddcd49bdc) = 0 -538166320188632:538166320188883 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291728112}, 1, 0x7ffddcd49be0) = 0 -538166320189474:538166320189724 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291728112}, 2, 0x7ffddcd49c00) = 0 -538166320190205:538166320190456 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291728112}, 6, 0x7ffddcd49c08) = 0 -538166320190205:538166320191347 882613:882613 hsa_amd_agent_iterate_memory_pools({handle=93863291726480}, 1, 0x7ffddcd49e18) = 0 -538166320194203:538166320196167 882613:882613 hsa_agent_get_info({handle=93863291737200}, 0, 0x7ffddcd49dd8) = 0 -538166320196698:538166320196958 882613:882613 hsa_agent_get_info({handle=93863291737200}, 17, 0x7ffddcd49dc0) = 0 -538166320198170:538166320198421 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291720576}, 0, 0x7ffddcd49bcc) = 0 -538166320198902:538166320199162 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291720576}, 1, 0x7ffddcd49bd0) = 0 -538166320199643:538166320199894 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291720576}, 2, 0x7ffddcd49bf0) = 0 -538166320200364:538166320200615 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291720576}, 6, 0x7ffddcd49bf8) = 0 -538166320201386:538166320201727 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291721456}, 0, 0x7ffddcd49bcc) = 0 -538166320201386:538166320202288 882613:882613 hsa_amd_agent_iterate_memory_pools({handle=93863291737200}, 1, 0x7ffddcd49e18) = 0 -538166320203881:538166320204332 882613:882613 hsa_agent_get_info({handle=93863291966720}, 0, 0x7ffddcd49dd8) = 0 -538166320204843:538166320205103 882613:882613 hsa_agent_get_info({handle=93863291966720}, 17, 0x7ffddcd49dc0) = 0 -538166320206366:538166320206606 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291968448}, 0, 0x7ffddcd49bcc) = 0 -538166320207087:538166320207348 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291968448}, 1, 0x7ffddcd49bd0) = 0 -538166320207818:538166320208069 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291968448}, 2, 0x7ffddcd49bf0) = 0 -538166320209782:538166320210043 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291968448}, 6, 0x7ffddcd49bf8) = 0 -538166320210724:538166320210984 882613:882613 hsa_amd_memory_pool_get_info({handle=93863291969440}, 0, 0x7ffddcd49bcc) = 0 -538166320210724:538166320211566 882613:882613 hsa_amd_agent_iterate_memory_pools({handle=93863291966720}, 1, 0x7ffddcd49e18) = 0 -538166320210724:538166320213028 882613:882613 hsa_iterate_agents(1, 0) = 0 -538166320294151:538166336777646 882613:882616 hsa_queue_create({handle=93863291737200}, 1024, 1, 0, 0, 0, 0, 0x7f91ed49ac88) = 0 -538166336809796:538166336869439 882613:882616 hsa_code_object_reader_create_from_file(5, 0x7f91ed49acb8) = 0 -538166336870721:538166336883686 882613:882616 hsa_executable_create_alt(1, 0, 0x0, 0x7f91ed49acc0) = 0 -538166336884868:538166337268842 882613:882616 hsa_executable_load_agent_code_object({handle=140259270599920}, {handle=93863291737200}, {handle=140259270599248}, 0x0, 0) = 0 -538166337270696:538166337974663 882613:882616 hsa_executable_freeze({handle=140259270599920}, 0x0) = 0 -538166337978250:538166337979312 882613:882616 hsa_executable_get_symbol_by_name({handle=140259270599920}, 0x7f90a80029b0, 0x7f91ed49abe8, 0x7f91ed49ac08) = 1013 -538166337982077:538166337982708 882613:882616 hsa_executable_get_symbol_by_name({handle=140259270599920}, 0x7f90a80038b0, 0x7f91ed49abe8, 0x7f91ed49ac08) = 0 -538166337984311:538166337984983 882613:882616 hsa_executable_symbol_get_info({handle=140259270612544}, 22, 0x7f91ed49acd0) = 0 -538166337985904:538166337986295 882613:882616 hsa_executable_symbol_get_info({handle=140259270612544}, 14, 0x7f91ed49acd8) = 0 -538166337987207:538166337987597 882613:882616 hsa_executable_symbol_get_info({handle=140259270612544}, 13, 0x7f91ed49acdc) = 0 -538166337988499:538166337988880 882613:882616 hsa_executable_symbol_get_info({handle=140259270612544}, 11, 0x7f91ed49ace0) = 0 -538166337989792:538166337990182 882613:882616 hsa_executable_symbol_get_info({handle=140259270612544}, 12, 0x7f91ed49ace4) = 0 -538166337991855:538166338004199 882613:882616 hsa_amd_memory_pool_allocate({handle=93863291727280}, 72, 0, 0x7f91ed49ac30) = 0 -538166338005371:538166338705110 882613:882616 hsa_amd_agents_allow_access(3, 0x555e40337c20, 0, 0x7f91ecc74000) = 0 -538166338706453:538166338719107 882613:882616 hsa_amd_memory_pool_allocate({handle=93863291727280}, 256, 0, 0x7f91ed49ac30) = 0 -538166338720059:538166338850665 882613:882616 hsa_amd_agents_allow_access(3, 0x555e40337c20, 0, 0x7f91ecc72000) = 0 -538166338851607:538166338864782 882613:882616 hsa_amd_memory_pool_allocate({handle=93863291727280}, 256, 0, 0x7f91ed49ac30) = 0 -538166338865723:538166338978927 882613:882616 hsa_amd_agents_allow_access(3, 0x555e40337c20, 0, 0x7f91ecc70000) = 0 -538166338980309:538166338982524 882613:882616 hsa_amd_signal_create(1, 0, 0, 0, 0x7f91ed49ac90) = 0 -538166338983816:538166338984628 882613:882616 hsa_queue_load_write_index_relaxed(0x7f91edff6000) = 0 -538166338985549:538166338986160 882613:882616 hsa_queue_load_read_index_relaxed(0x7f91edff6000) = 0 -538166338987032:538166338987453 882613:882616 hsa_queue_store_write_index_screlease(0x7f91edff6000, 1) = void -538166338988515:538166338989397 882613:882616 hsa_signal_store_screlease({handle=140264755457280}, 0) = void -538166338990278:538166339006328 882613:882616 hsa_signal_wait_scacquire({handle=140264755456768}, 0, 0, 18446744073709551615, 0) = 0 -538166339009224:538166339010346 882613:882616 hsa_signal_destroy({handle=140264755456768}) = 0 -538166339011689:538166339037938 882613:882616 hsa_memory_free(0x7f91ecc72000) = 0 -538166339038930:538166339052135 882613:882616 hsa_memory_free(0x7f91ecc70000) = 0 -538166339053307:538166339082843 882613:882616 hsa_executable_destroy({handle=140259270599920}) = 0 -538166339083895:538166339087181 882613:882616 hsa_code_object_reader_destroy({handle=140259270599248}) = 0 +975779240024464:975779240024815 566828:566828 hsa_agent_get_info(, 0, 0x7ffdd25cc9a8) = 0 :6 +975779240029173:975779240029274 566828:566828 hsa_agent_get_info(, 17, 0x7ffdd25cc990) = 0 :7 +975779240035816:975779240036187 566828:566828 hsa_amd_memory_pool_get_info(, 0, 0x7ffdd25cc6bc) = 0 :9 +975779240036667:975779240036768 566828:566828 hsa_amd_memory_pool_get_info(, 1, 0x7ffdd25cc6c0) = 0 :10 +975779240037219:975779240037319 566828:566828 hsa_amd_memory_pool_get_info(, 2, 0x7ffdd25cc6e0) = 0 :11 +975779240037760:975779240037860 566828:566828 hsa_amd_memory_pool_get_info(, 6, 0x7ffdd25cc6e8) = 0 :12 +975779240039823:975779240039914 566828:566828 hsa_amd_memory_pool_get_info(, 0, 0x7ffdd25cc6bc) = 0 :13 +975779240040455:975779240040555 566828:566828 hsa_amd_memory_pool_get_info(, 1, 0x7ffdd25cc6c0) = 0 :14 +975779240041076:975779240041156 566828:566828 hsa_amd_memory_pool_get_info(, 2, 0x7ffdd25cc6e0) = 0 :15 +975779240041697:975779240041777 566828:566828 hsa_amd_memory_pool_get_info(, 6, 0x7ffdd25cc6e8) = 0 :16 +975779240042619:975779240042709 566828:566828 hsa_amd_memory_pool_get_info(, 0, 0x7ffdd25cc6bc) = 0 :17 +975779240043250:975779240043350 566828:566828 hsa_amd_memory_pool_get_info(, 1, 0x7ffdd25cc6c0) = 0 :18 +975779240043871:975779240043961 566828:566828 hsa_amd_memory_pool_get_info(, 2, 0x7ffdd25cc6e0) = 0 :19 +975779240044482:975779240044562 566828:566828 hsa_amd_memory_pool_get_info(, 6, 0x7ffdd25cc6e8) = 0 :20 +975779240044482:975779240045264 566828:566828 hsa_amd_agent_iterate_memory_pools(, 1, 0x7ffdd25cc9e8) = 0 :8 +975779240048430:975779240049341 566828:566828 hsa_agent_get_info(, 0, 0x7ffdd25cc9a8) = 0 :21 +975779240049822:975779240049932 566828:566828 hsa_agent_get_info(, 17, 0x7ffdd25cc990) = 0 :22 +975779240050654:975779240050744 566828:566828 hsa_amd_memory_pool_get_info(, 0, 0x7ffdd25cc62c) = 0 :24 +975779240051275:975779240051365 566828:566828 hsa_amd_memory_pool_get_info(, 1, 0x7ffdd25cc630) = 0 :25 +975779240051786:975779240051866 566828:566828 hsa_amd_memory_pool_get_info(, 2, 0x7ffdd25cc650) = 0 :26 +975779240052287:975779240052377 566828:566828 hsa_amd_memory_pool_get_info(, 6, 0x7ffdd25cc658) = 0 :27 +975779240053048:975779240053159 566828:566828 hsa_amd_memory_pool_get_info(, 0, 0x7ffdd25cc62c) = 0 :28 +975779240053048:975779240053609 566828:566828 hsa_amd_agent_iterate_memory_pools(, 1, 0x7ffdd25cc9e8) = 0 :23 +975779240055373:975779240055663 566828:566828 hsa_agent_get_info(, 0, 0x7ffdd25cc9a8) = 0 :29 +975779240056144:975779240056234 566828:566828 hsa_agent_get_info(, 17, 0x7ffdd25cc990) = 0 :30 +975779240056986:975779240057076 566828:566828 hsa_amd_memory_pool_get_info(, 0, 0x7ffdd25cc62c) = 0 :32 +975779240057507:975779240057587 566828:566828 hsa_amd_memory_pool_get_info(, 1, 0x7ffdd25cc630) = 0 :33 +975779240058008:975779240058088 566828:566828 hsa_amd_memory_pool_get_info(, 2, 0x7ffdd25cc650) = 0 :34 +975779240058509:975779240058589 566828:566828 hsa_amd_memory_pool_get_info(, 6, 0x7ffdd25cc658) = 0 :35 +975779240061504:975779240061605 566828:566828 hsa_amd_memory_pool_get_info(, 0, 0x7ffdd25cc62c) = 0 :36 +975779240061504:975779240062035 566828:566828 hsa_amd_agent_iterate_memory_pools(, 1, 0x7ffdd25cc9e8) = 0 :31 +975779240061504:975779240063528 566828:566828 hsa_iterate_agents(1, 0) = 0 :5 +975779240167184:975779249865797 566828:566834 hsa_queue_create(, 1024, 1, 0, 0, 0, 0, 0x7f70535fdbc8) = 0 :37 +975779249901595:975779249990022 566828:566834 hsa_code_object_reader_create_from_file(8, 0x7f70535fdbf8) = 0 :38 +975779249990984:975779250001383 566828:566834 hsa_executable_create_alt(1, 0, 0x0, 0x7f70535fdc00) = 0 :27 +975779250002345:975779250430202 566828:566834 hsa_executable_load_agent_code_object(, , , 0x0, 0) = 0 :28 +975779250432296:975779250926909 566828:566834 hsa_executable_freeze(, 0x0) = 0 :29 +975779250929755:975779250931207 566828:566834 hsa_executable_get_symbol_by_name(, 0x7f7054002950, 0x7f70535fdb28, 0x7f70535fdb48) = 1013 :2a +975779250931979:975779250932690 566828:566834 hsa_executable_get_symbol_by_name(, 0x7f7054003850, 0x7f70535fdb28, 0x7f70535fdb48) = 0 :2b +975779250934403:975779250934854 566828:566834 hsa_executable_symbol_get_info(, 22, 0x7f70535fdc10) = 0 :44 +975779250935425:975779250935536 566828:566834 hsa_executable_symbol_get_info(, 14, 0x7f70535fdc18) = 0 :45 +975779250936097:975779250936177 566828:566834 hsa_executable_symbol_get_info(, 13, 0x7f70535fdc1c) = 0 :46 +975779250936728:975779250936798 566828:566834 hsa_executable_symbol_get_info(, 11, 0x7f70535fdc20) = 0 :47 +975779250937349:975779250937419 566828:566834 hsa_executable_symbol_get_info(, 12, 0x7f70535fdc24) = 0 :48 +975779250938321:975779250956876 566828:566834 hsa_amd_memory_pool_allocate(, 72, 0, 0x7f70535fdb70) = 0 :49 +975779250958098:975779251048298 566828:566834 hsa_amd_agents_allow_access(3, 0x55e1b9df4c30, 0, 0x7f7261070000) = 0 :50 +975779251049150:975779251065531 566828:566834 hsa_amd_memory_pool_allocate(, 256, 0, 0x7f70535fdb70) = 0 :51 +975779251066232:975779251149319 566828:566834 hsa_amd_agents_allow_access(3, 0x55e1b9df4c30, 0, 0x7f726106e000) = 0 :52 +975779251150000:975779251165960 566828:566834 hsa_amd_memory_pool_allocate(, 256, 0, 0x7f70535fdb70) = 0 :53 +975779251166531:975779251256912 566828:566834 hsa_amd_agents_allow_access(3, 0x55e1b9df4c30, 0, 0x7f726106c000) = 0 :54 +975779251258114:975779251261601 566828:566834 hsa_amd_signal_create(1, 0, 0, 0, 0x7f70535fdbd0) = 0 :55 +975779251262923:975779251263204 566828:566834 hsa_queue_load_write_index_relaxed(0x7f726109e000) = 0 :56 +975779251264065:975779251264276 566828:566834 hsa_queue_load_read_index_relaxed(0x7f726109e000) = 0 :57 +975779251264937:975779251265178 566828:566834 hsa_queue_store_write_index_screlease(0x7f726109e000, 1) = void :58 +975779251265969:975779251266951 566828:566834 hsa_signal_store_screlease(, 0) = void :59 +975779251267472:975779251283773 566828:566834 hsa_signal_wait_scacquire(, 0, 0, 18446744073709551615, 0) = 0 :60 +975779251284654:975779251286848 566828:566834 hsa_signal_destroy() = 0 :61 +975779251290806:975779251322035 566828:566834 hsa_memory_free(0x7f726106e000) = 0 :62 +975779251322646:975779251341261 566828:566834 hsa_memory_free(0x7f726106c000) = 0 :63 +975779251342043:975779251389061 566828:566834 hsa_executable_destroy() = 0 :64 +975779251389843:975779251392488 566828:566834 hsa_code_object_reader_destroy() = 0 :65 diff --git a/test/golden_traces/load_unload_reload_trace.txt b/test/golden_traces/load_unload_reload_trace.txt index a902501de5..2cc1e01e7e 100644 --- a/test/golden_traces/load_unload_reload_trace.txt +++ b/test/golden_traces/load_unload_reload_trace.txt @@ -1,24 +1,22 @@ -ROCTracer (pid=202688): -0x55a6e1bfa280 agent cpu -0x55a6e1bf9470 agent gpu -0x55a6e1c34d10 agent gpu -169339419628779 +ROCTracer (pid=566858): +0x55ae2fa607c0 agent cpu +0x55ae2fa5eeb0 agent gpu +0x55ae2fa9b540 agent gpu +975785718853775 HSA-trace() HSA-activity-trace() -169339419776888:169339419781697 202688:202688 hsa_amd_profiling_async_copy_enable(1) = 0 -169339420102142:169339420102703 202688:202688 hsa_agent_get_info({handle=94175240364672}, 17, 0x7ffe818aff34) = 0 -169339420104075:169339420104606 202688:202688 hsa_agent_get_info({handle=94175240361072}, 17, 0x7ffe818aff34) = 0 -169339420105568:169339420106049 202688:202688 hsa_agent_get_info({handle=94175240604944}, 17, 0x7ffe818aff34) = 0 -169339420105568:169339420106941 202688:202688 hsa_iterate_agents(1, 0) = 0 -ROCTracer (pid=202688): -0x55a6e1bfa280 agent cpu -0x55a6e1c99260 agent gpu -0x55a6e1c95d50 agent gpu -169339843452619 +975785719398623:975785719398824 566858:566858 hsa_agent_get_info(, 17, 0x7ffe30b97814) = 0 :6 +975785719403482:975785719403643 566858:566858 hsa_agent_get_info(, 17, 0x7ffe30b97814) = 0 :7 +975785719404274:975785719404364 566858:566858 hsa_agent_get_info(, 17, 0x7ffe30b97814) = 0 :8 +975785719404274:975785719404885 566858:566858 hsa_iterate_agents(1, 0) = 0 :5 +ROCTracer (pid=566858): +0x55ae2fa607c0 agent cpu +0x55ae2fb02cc0 agent gpu +0x55ae2fa62970 agent gpu +975785742239830 HSA-trace() HSA-activity-trace() -169339843601029:169339843605668 202688:202688 hsa_amd_profiling_async_copy_enable(1) = 0 -169339843768255:169339843768816 202688:202688 hsa_agent_get_info({handle=94175240364672}, 17, 0x7ffe818aff34) = 0 -169339843770028:169339843770549 202688:202688 hsa_agent_get_info({handle=94175241015904}, 17, 0x7ffe818aff34) = 0 -169339843771491:169339843771962 202688:202688 hsa_agent_get_info({handle=94175241002320}, 17, 0x7ffe818aff34) = 0 -169339843771491:169339843772854 202688:202688 hsa_iterate_agents(1, 0) = 0 +975785742436120:975785742436310 566858:566858 hsa_agent_get_info(, 17, 0x7ffe30b97814) = 0 :14 +975785742437352:975785742437472 566858:566858 hsa_agent_get_info(, 17, 0x7ffe30b97814) = 0 :15 +975785742437963:975785742438053 566858:566858 hsa_agent_get_info(, 17, 0x7ffe30b97814) = 0 :16 +975785742437963:975785742438464 566858:566858 hsa_iterate_agents(1, 0) = 0 :13