diff --git a/projects/rocprofiler/src/api/rocmtool.cpp b/projects/rocprofiler/src/api/rocmtool.cpp index 3015be744e..fb1a5405af 100644 --- a/projects/rocprofiler/src/api/rocmtool.cpp +++ b/projects/rocprofiler/src/api/rocmtool.cpp @@ -199,7 +199,7 @@ size_t rocmtool::GetKernelInfoSize(rocprofiler_kernel_info_kind_t kind, rocprofiler_kernel_id_t kernel_id) { switch (kind) { case ROCPROFILER_KERNEL_NAME: - return GetKernelNameFromKsymbols(kernel_id.handle).size(); + return GetKernelNameUsingDispatchID(kernel_id.handle).size(); default: warning("The provided Kernel Kind is not yet supported!"); return 0; @@ -209,7 +209,7 @@ const char* rocmtool::GetKernelInfo(rocprofiler_kernel_info_kind_t kind, rocprofiler_kernel_id_t kernel_id) { switch (kind) { case ROCPROFILER_KERNEL_NAME: - return strdup(GetKernelNameFromKsymbols(kernel_id.handle).c_str()); + return strdup(GetKernelNameUsingDispatchID(kernel_id.handle).c_str()); default: warning("The provided Kernel Kind is not yet supported!"); return ""; diff --git a/projects/rocprofiler/src/core/hsa/hsa_support.cpp b/projects/rocprofiler/src/core/hsa/hsa_support.cpp index b2448cc415..c3c22abf7b 100644 --- a/projects/rocprofiler/src/core/hsa/hsa_support.cpp +++ b/projects/rocprofiler/src/core/hsa/hsa_support.cpp @@ -52,6 +52,7 @@ namespace { hsa_status_t hsa_executable_iteration_callback(hsa_executable_t executable, hsa_agent_t agent, hsa_executable_symbol_t symbol, void* args) { + hsa_symbol_kind_t type; rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn( symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &type); @@ -62,14 +63,21 @@ hsa_status_t hsa_executable_iteration_callback(hsa_executable_t executable, hsa_ // TODO(aelwazir): to be removed if the HSA fixed the issue of corrupted // names overflowing the length given if (name_length > 1) { - char name[name_length + 1]; - uint64_t kernel_object; - rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn( - symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name); - rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn( - symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object); - std::string kernel_name = std::string(name).substr(0, name_length); - rocmtools::AddKernelName(kernel_object, kernel_name); + if(!(*static_cast(args))) { + char name[name_length + 1]; + uint64_t kernel_object; + rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn( + symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name); + rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn( + symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object); + std::string kernel_name = std::string(name).substr(0, name_length); + rocmtools::AddKernelName(kernel_object, kernel_name); + } else { + uint64_t kernel_object; + rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn( + symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object); + rocmtools::RemoveKernelName(kernel_object); + } } } @@ -447,7 +455,7 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable, ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, &data); hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent, - hsa_executable_iteration_callback, nullptr); + hsa_executable_iteration_callback, &(data.codeobj.unload)); return HSA_STATUS_SUCCESS; } diff --git a/projects/rocprofiler/src/core/hsa/queues/queue.cpp b/projects/rocprofiler/src/core/hsa/queues/queue.cpp index cb51fa35ce..ea930d2b3b 100644 --- a/projects/rocprofiler/src/core/hsa/queues/queue.cpp +++ b/projects/rocprofiler/src/core/hsa/queues/queue.cpp @@ -63,6 +63,10 @@ void AddKernelName(uint64_t handle, std::string name) { std::lock_guard lock(ksymbol_map_lock); ksymbols->emplace(handle, name); } +void RemoveKernelName(uint64_t handle) { + std::lock_guard lock(ksymbol_map_lock); + ksymbols->erase(handle); +} std::string GetKernelNameFromKsymbols(uint64_t handle) { std::lock_guard lock(ksymbol_map_lock); if(ksymbols->find(handle)!=ksymbols->end()) @@ -70,11 +74,39 @@ std::string GetKernelNameFromKsymbols(uint64_t handle) { else return "Unknown Kernel!"; } + +static std::mutex kernel_names_map_lock; +static std::map>* kernel_names; +static std::atomic kernel_names_flag{true}; +void AddKernelNameWithDispatchID(std::string name, uint64_t id) { + std::lock_guard lock(kernel_names_map_lock); + if(kernel_names->find(name) == kernel_names->end()) + kernel_names->emplace(name, std::vector()); + kernel_names->at(name).push_back(id); +} +std::string GetKernelNameUsingDispatchID(uint64_t given_id) { + std::lock_guard lock(kernel_names_map_lock); + for(auto kernel_name : (*kernel_names)) { + for(auto dispatch_id : kernel_name.second) { + if(dispatch_id == given_id) + return kernel_name.first; + } + } + return "Unknown Kernel!"; +} + void InitKsymbols() { if (ksymbols_flag.load(std::memory_order_relaxed)) { - std::lock_guard lock(ksymbol_map_lock); - ksymbols = new std::map(); - ksymbols_flag.exchange(false, std::memory_order_release); + { + std::lock_guard lock(ksymbol_map_lock); + ksymbols = new std::map(); + ksymbols_flag.exchange(false, std::memory_order_release); + } + { + std::lock_guard lock(kernel_names_map_lock); + kernel_names = new std::map>(); + kernel_names_flag.exchange(false, std::memory_order_release); + } } } void FinitKsymbols() { @@ -84,8 +116,16 @@ void FinitKsymbols() { delete ksymbols; ksymbols_flag.exchange(true, std::memory_order_release); } + if (!kernel_names_flag.load(std::memory_order_relaxed)) { + std::lock_guard lock(kernel_names_map_lock); + kernel_names->clear(); + delete kernel_names; + kernel_names_flag.exchange(true, std::memory_order_release); + } } + + struct kernel_descriptor_t { uint8_t reserved0[16]; int64_t kernel_code_entry_byte_offset; @@ -416,7 +456,6 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { hsa_support::GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn( queue_info_session->agent, pending.signal, &time); rocprofiler_record_profiler_t record{}; - record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor}; record.gpu_id = rocprofiler_agent_id_t{ (uint64_t)hsa_support::GetAgentInfo(queue_info_session->agent.handle).getIndex()}; record.kernel_properties = pending.kernel_properties; @@ -429,7 +468,8 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { AddRecordCounters(&record, pending); } record.header = {ROCPROFILER_PROFILER_RECORD, - rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}}; + rocprofiler_record_id_t{pending.kernel_descriptor}}; + record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor}; if (pending.session_id.handle == 0) { pending.session_id = GetROCMToolObj()->GetCurrentSessionId(); @@ -509,7 +549,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { AddAttRecord(&record, queue_info_session->agent, pending); } record.header = {ROCPROFILER_ATT_TRACER_RECORD, - rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}}; + rocprofiler_record_id_t{pending.kernel_descriptor}}; if (pending.session_id.handle == 0) { pending.session_id = GetROCMToolObj()->GetCurrentSessionId(); @@ -740,14 +780,16 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt rocprofiler_kernel_properties_t kernel_properties = set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent()); if (session) { + uint64_t record_id = GetROCMToolObj()->GetUniqueRecordId(); + AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id); if (profiles && replay_mode_count > 0) { session->GetProfiler()->AddPendingSignals( - writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, + writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, profile.first, profile.first->metrics_list.size(), profile.second, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); } else { session->GetProfiler()->AddPendingSignals( - writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, + writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, nullptr, 0, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); } @@ -929,13 +971,15 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt // list to be processed by the signal interrupt rocprofiler_kernel_properties_t kernel_properties = set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent()); + uint64_t record_id = GetROCMToolObj()->GetUniqueRecordId(); + AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id); if (session && profile) { session->GetAttTracer()->AddPendingSignals( - writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id, + writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, profile, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); } else { session->GetAttTracer()->AddPendingSignals( - writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id, + writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); } diff --git a/projects/rocprofiler/src/core/hsa/queues/queue.h b/projects/rocprofiler/src/core/hsa/queues/queue.h index 1999b1ce2c..5ed0ae798b 100644 --- a/projects/rocprofiler/src/core/hsa/queues/queue.h +++ b/projects/rocprofiler/src/core/hsa/queues/queue.h @@ -42,6 +42,9 @@ namespace rocmtools { void InitKsymbols(); void FinitKsymbols(); void AddKernelName(uint64_t handle, std::string kernel_name); +void RemoveKernelName(uint64_t handle); +void AddKernelNameWithDispatchID(std::string name, uint64_t id); +std::string GetKernelNameUsingDispatchID(uint64_t given_id); std::string GetKernelNameFromKsymbols(uint64_t handle); uint32_t GetCurrentActiveInterruptSignalsCount();