From 41be168eb33bf96f61052c87fe20800d725aa532 Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Thu, 1 Jun 2023 21:26:20 +0000 Subject: [PATCH] Fixing Kernel Dispatches timestamps issue Change-Id: I4946ea13e2796b018dc3a9114c232b0fa5f483dc [ROCm/rocprofiler commit: 941a0373e6c7f402497e9c27586a7168635210f9] --- .../rocprofiler/src/core/hsa/queues/queue.cpp | 25 +++++++------------ .../rocprofiler/src/core/session/att/att.cpp | 4 +-- .../rocprofiler/src/core/session/att/att.h | 5 ++-- .../src/core/session/profiler/profiler.cpp | 4 +-- .../src/core/session/profiler/profiler.h | 5 ++-- 5 files changed, 19 insertions(+), 24 deletions(-) diff --git a/projects/rocprofiler/src/core/hsa/queues/queue.cpp b/projects/rocprofiler/src/core/hsa/queues/queue.cpp index 12da0ca807..cbdc14dff6 100644 --- a/projects/rocprofiler/src/core/hsa/queues/queue.cpp +++ b/projects/rocprofiler/src/core/hsa/queues/queue.cpp @@ -397,10 +397,10 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { for (auto it = pending_signals.begin(); it != pending_signals.end(); it = pending_signals.erase(it)) { auto& pending = *it; - if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->signal)) return true; + if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal)) return true; hsa_amd_profiling_dispatch_time_t time; hsa_support::GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn( - queue_info_session->agent, pending->signal, &time); + queue_info_session->agent, pending->original_signal, &time); uint32_t record_count = 1; bool is_individual_xcc_mode = false; uint32_t xcc_count = queue_info_session->xcc_count; @@ -469,8 +469,8 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { } delete pending->context; } - if (pending->signal.handle) - hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(pending->signal); + if (pending->new_signal.handle) + hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal); if (queue_info_session->interrupt_signal.handle) hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(queue_info_session->interrupt_signal); } @@ -500,7 +500,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { it = pending_signals.erase(it)) { auto& pending = *it; std::lock_guard lock(session->GetSessionLock()); - if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.signal)) return true; + if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.new_signal)) return true; rocprofiler_record_att_tracer_t record{}; record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor}; record.gpu_id = rocprofiler_agent_id_t{(uint64_t)queue_info_session->gpu_index}; @@ -725,13 +725,6 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt auto& packet = transformed_packets.emplace_back(packets_arr[i]); auto& dispatch_packet = reinterpret_cast(packet); - - /* - * Only PC sampling relies on this right now, so it would be better to - * only generate an ID if PC sampling is active to conserve IDs, but it's - * unlikely 64 bits' worth of identifiers will be exhausted during the - * lifetime of the ROCProfiler_Singleton. - */ uint64_t correlation_id = dispatch_packet.reserved2; CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal); @@ -745,12 +738,12 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt record_id); if (profiles.size() > 0 && replay_mode_count > 0) { session->GetProfiler()->AddPendingSignals( - writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, + writer_id, record_id, original_packet.completion_signal, 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, correlation_id); } else { session->GetProfiler()->AddPendingSignals( - writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, + writer_id, record_id, original_packet.completion_signal, dispatch_packet.completion_signal, session_id, buffer_id, nullptr, 0, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id); } @@ -946,11 +939,11 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt record_id); if (session && profile) { session->GetAttTracer()->AddPendingSignals( - writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, profile, + writer_id, record_id, original_packet.completion_signal, 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, record_id, dispatch_packet.completion_signal, session_id, buffer_id, nullptr, + writer_id, record_id, original_packet.completion_signal, 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/session/att/att.cpp b/projects/rocprofiler/src/core/session/att/att.cpp index e824fdac2c..d309022c2c 100644 --- a/projects/rocprofiler/src/core/session/att/att.cpp +++ b/projects/rocprofiler/src/core/session/att/att.cpp @@ -31,7 +31,7 @@ AttTracer::AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t AttTracer::~AttTracer() {} void AttTracer::AddPendingSignals(uint32_t writer_id, uint64_t kernel_object, - const hsa_signal_t& completion_signal, + const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id, hsa_ven_amd_aqlprofile_profile_t* profile, @@ -41,7 +41,7 @@ void AttTracer::AddPendingSignals(uint32_t writer_id, uint64_t kernel_object, if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end()) sessions_pending_signals_.emplace(writer_id, std::vector()); sessions_pending_signals_.at(writer_id).emplace_back( - att_pending_signal_t{kernel_object, completion_signal, session_id_, buffer_id, profile, + att_pending_signal_t{kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, profile, kernel_properties, thread_id, queue_index}); } diff --git a/projects/rocprofiler/src/core/session/att/att.h b/projects/rocprofiler/src/core/session/att/att.h index 19214577ef..9950b3bcbe 100644 --- a/projects/rocprofiler/src/core/session/att/att.h +++ b/projects/rocprofiler/src/core/session/att/att.h @@ -34,7 +34,8 @@ namespace rocprofiler { typedef struct { uint64_t kernel_descriptor; - hsa_signal_t signal; + hsa_signal_t original_signal; + hsa_signal_t new_signal; rocprofiler_session_id_t session_id; rocprofiler_buffer_id_t buffer_id; hsa_ven_amd_aqlprofile_profile_t* profile; @@ -52,7 +53,7 @@ class AttTracer { ~AttTracer(); void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object, - const hsa_signal_t& completion_signal, rocprofiler_session_id_t session_id, + const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id, hsa_ven_amd_aqlprofile_profile_t* profile, rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, uint64_t queue_index); diff --git a/projects/rocprofiler/src/core/session/profiler/profiler.cpp b/projects/rocprofiler/src/core/session/profiler/profiler.cpp index 7baacb1963..c7828066b7 100644 --- a/projects/rocprofiler/src/core/session/profiler/profiler.cpp +++ b/projects/rocprofiler/src/core/session/profiler/profiler.cpp @@ -125,7 +125,7 @@ bool Profiler::HasActivePass() { } void Profiler::AddPendingSignals( - uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& completion_signal, + uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id, rocprofiler::profiling_context_t* context, uint64_t session_data_count, hsa_ven_amd_aqlprofile_profile_t* profile, rocprofiler_kernel_properties_t kernel_properties, @@ -134,7 +134,7 @@ void Profiler::AddPendingSignals( if (sessions_pending_signals_->find(writer_id) == sessions_pending_signals_->end()) sessions_pending_signals_->emplace(writer_id, std::vector()); sessions_pending_signals_->at(writer_id).emplace_back(new pending_signal_t{ - kernel_object, completion_signal, session_id_, buffer_id, context, session_data_count, + kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, context, session_data_count, profile, kernel_properties, thread_id, queue_index, correlation_id}); } diff --git a/projects/rocprofiler/src/core/session/profiler/profiler.h b/projects/rocprofiler/src/core/session/profiler/profiler.h index b61b210457..5e06fb1c22 100644 --- a/projects/rocprofiler/src/core/session/profiler/profiler.h +++ b/projects/rocprofiler/src/core/session/profiler/profiler.h @@ -44,7 +44,8 @@ namespace rocprofiler { typedef struct { uint64_t kernel_descriptor; - hsa_signal_t signal; + hsa_signal_t original_signal; + hsa_signal_t new_signal; rocprofiler_session_id_t session_id; rocprofiler_buffer_id_t buffer_id; rocprofiler::profiling_context_t* context; @@ -67,7 +68,7 @@ class Profiler { ~Profiler(); void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object, - const hsa_signal_t& completion_signal, rocprofiler_session_id_t session_id, + const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id, rocprofiler::profiling_context_t* context, uint64_t session_data_count, hsa_ven_amd_aqlprofile_profile_t* profile,