From be053cb10bfd92339d03b94c03d6d27a8deeb98c Mon Sep 17 00:00:00 2001 From: "Trowbridge, Ian" Date: Thu, 13 Mar 2025 22:40:03 -0500 Subject: [PATCH] Temporarily Fix Incorrect Kernel Perfetto Trace Duration due to Firmware Timestamp Bug (#134) * Perfetto duration temp fix setup * Add timestamp change amounts to ROCP Info * Groups kernel dispatch info by agent and queue id before sorting. Midpoint interpolation is then performed on the sorted kernels * Moved dispatch bins into the for-loop * Fix compilation error by using const ref * Modified for review comments * Changed variable names [ROCm/rocprofiler-sdk commit: 6518c5463dc71477c3119c7bb26770c906839428] --- .../source/lib/output/generatePerfetto.cpp | 162 ++++++++++++------ 1 file changed, 110 insertions(+), 52 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp b/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp index f9886fda3f..657ba3bb92 100644 --- a/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp +++ b/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp @@ -32,8 +32,11 @@ #include #include +#include + #include #include +#include #include #include #include @@ -505,62 +508,117 @@ write_perfetto( itr.end_timestamp); tracing_session->FlushBlocking(); } - for(auto ditr : kernel_dispatch_gen) - for(auto itr : kernel_dispatch_gen.get(ditr)) + { + auto generator = kernel_dispatch_gen.get(ditr); + // Group kernels on the same queue and agent. Temporary fix for firmware timestamp bug + // Can be removed once bug is resolved. + auto dispatch_bins = std::unordered_map< + rocprofiler_agent_id_t, + std::unordered_map< + rocprofiler_queue_id_t, + std::vector>>{}; + for(auto& itr : generator) { - const auto& info = itr.dispatch_info; - const kernel_symbol_info* sym = tool_metadata.get_kernel_symbol(info.kernel_id); - - CHECK(sym != nullptr); - - auto name = std::string_view{sym->kernel_name}; - auto& track = agent_queue_tracks.at(info.agent_id).at(info.queue_id); - - if(demangled.find(name) == demangled.end()) - { - demangled.emplace(name, common::cxx_demangle(name)); - } - - TRACE_EVENT_BEGIN( - sdk::perfetto_category::name, - ::perfetto::StaticString(demangled.at(name).c_str()), - track, - itr.start_timestamp, - ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), - "begin_ns", - itr.start_timestamp, - "end_ns", - itr.end_timestamp, - "delta_ns", - (itr.end_timestamp - itr.start_timestamp), - "kind", - itr.kind, - "agent", - agents_map.at(info.agent_id).logical_node_id, - "corr_id", - itr.correlation_id.internal, - "queue", - info.queue_id.handle, - "tid", - itr.thread_id, - "kernel_id", - info.kernel_id, - "private_segment_size", - info.private_segment_size, - "group_segment_size", - info.group_segment_size, - "workgroup_size", - info.workgroup_size.x * info.workgroup_size.y * info.workgroup_size.z, - "grid_size", - info.grid_size.x * info.grid_size.y * info.grid_size.z); - TRACE_EVENT_END(sdk::perfetto_category::name, - track, - itr.end_timestamp); - tracing_session->FlushBlocking(); + const auto& info = itr.dispatch_info; + dispatch_bins[info.agent_id][info.queue_id].emplace_back(&itr); } - } + for(const auto& aitr : dispatch_bins) + { + for(auto qitr : aitr.second) + { + // Sort kernels on the same queue and agent by timestamp + std::sort(qitr.second.begin(), + qitr.second.end(), + [](const auto* lhs, const auto* rhs) { + return lhs->start_timestamp < rhs->start_timestamp; + }); + + // Loop over the kernels (qitr.second) and put them into perfetto. + for(auto it = qitr.second.begin(); it != qitr.second.end(); ++it) + { + auto& current = **it; + const auto& info = current.dispatch_info; + const kernel_symbol_info* sym = + tool_metadata.get_kernel_symbol(info.kernel_id); + + CHECK(sym != nullptr); + + auto name = std::string_view{sym->kernel_name}; + auto& track = agent_queue_tracks.at(info.agent_id).at(info.queue_id); + + // Temporary fix until timestamp issues are resolved: Set timestamps to be + // halfway between ending timestamp and starting timestamp of overlapping + // kernel dispatches. Perfetto displays slices incorrectly if overlapping + // slices on the same track are not completely enveloped. + auto next = std::next(it); + if(next != qitr.second.end() && + (*next)->start_timestamp < (*it)->end_timestamp) + { + auto start = (*next)->start_timestamp; + auto end = std::min((*it)->end_timestamp, (*next)->end_timestamp); + auto mid = start + (end - start) / 2; + // Report changed timestamps to ROCP INFO + ROCP_INFO << fmt::format( + "Kernel ending timestamp increased by {} ns to {} ns with " + "following kernel starting timestamp decreased by {} ns to {} ns " + "due to firmware timestamp error.", + ((*it)->end_timestamp - mid), + mid, + (mid - (*next)->start_timestamp), + mid); + (*it)->end_timestamp = mid; + (*next)->start_timestamp = mid; + } + + if(demangled.find(name) == demangled.end()) + { + demangled.emplace(name, common::cxx_demangle(name)); + } + + TRACE_EVENT_BEGIN( + sdk::perfetto_category::name, + ::perfetto::StaticString(demangled.at(name).c_str()), + track, + current.start_timestamp, + ::perfetto::Flow::ProcessScoped(current.correlation_id.internal), + "begin_ns", + current.start_timestamp, + "end_ns", + current.end_timestamp, + "delta_ns", + (current.end_timestamp - current.start_timestamp), + "kind", + current.kind, + "agent", + tool_metadata.get_node_id(info.agent_id), + "corr_id", + current.correlation_id.internal, + "queue", + info.queue_id.handle, + "tid", + current.thread_id, + "kernel_id", + info.kernel_id, + "private_segment_size", + info.private_segment_size, + "group_segment_size", + info.group_segment_size, + "workgroup_size", + info.workgroup_size.x * info.workgroup_size.y * info.workgroup_size.z, + "grid_size", + info.grid_size.x * info.grid_size.y * info.grid_size.z); + TRACE_EVENT_END( + sdk::perfetto_category::name, + track, + current.end_timestamp); + tracing_session->FlushBlocking(); + } + } + } + } + } // counter tracks { // memory copy counter track