From d34d442d52d840c3e44ae3ef5cad8972ee617aee Mon Sep 17 00:00:00 2001 From: ajassani Date: Sun, 24 May 2026 03:11:13 +0000 Subject: [PATCH 1/3] rocprofiler-sdk-tool: surface inter-stream event dependencies on HIP wait/sync records rocprofv3 already shows per-record stream IDs on hipStreamWaitEvent and hipEventSynchronize tool records, but it does not yet say *which* stream and *which* hipEventRecord the waiter is actually depending on. This makes it hard for users to reconstruct critical-path / cross-stream dependency graphs from a captured trace without re-running the workload under a separate analysis pass. This change adds an event-lifecycle tracking layer inside the tool library and surfaces the resolved producer on the waiter's tool record in three of the existing output formats (Perfetto, JSON, CSV). OTF2 and rocpd are intentionally left for a follow-up PR because both involve schema-level changes (pre-declared OTF2 attribute IDs, SQL migration) that deserve separate review. Mechanism ========= A new utility in source/lib/rocprofiler-sdk-tool/event_producer_map.{hpp,cpp} maintains two process-global, shared/unique-locked maps using the existing common::Synchronized + common::static_object building blocks: * producer map: hipEvent_t -> { stream_id, hipEventRecord corr_id } written on PHASE_EXIT of hipEventRecord (+ _spt + hipEventRecordWithFlags), erased on hipEventDestroy * resolved-wait map: waiter corr_id -> { wait_on_stream, wait_on_corr_id } written on PHASE_EXIT of hipStreamWaitEvent (+ _spt) and hipEventSynchronize after a producer lookup, drained at buffer-flush time A dedicated callback context is registered on ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API filtered to just the seven operation IDs involved, gated on hip_runtime_api_trace so there is no overhead when HIP tracing is off. The resolved-wait map is necessary because the waiter's tool record is materialized at buffer-flush time -- a different code path from where the typed wait/sync args are accessible. Bridging the two via the waiter's correlation_id keeps the producer lookup at callback time (when the args are live) and the record population at flush time (when the tool record struct is being built). Tool record extension ===================== tool_buffer_tracing_hip_api_ext_record_t (lib/output/stream_info.hpp) gains two fields with backward-compatible defaults: rocprofiler_stream_id_t wait_on_stream = {}; // 0 = none uint64_t wait_on_correlation_id = 0; // 0 = none A 4-arg constructor variant accepts the resolved info; the existing 2-arg constructor is unchanged so non-wait HIP API records continue to work without modification. The cereal save() helper picks up the new fields so JSON output gets them for free. Output formats ============== * Perfetto: hipStreamWaitEvent / hipEventSynchronize slices gain `wait_on_stream_ID` and `wait_on_corr_id` arg fields and, when wait_on_correlation_id is non-zero, a second Flow::ProcessScoped(wait_on_correlation_id) so the dependency renders as an arrow from the producing hipEventRecord call to the waiter. * JSON: every hip_api buffer record now carries wait_on_stream and wait_on_correlation_id alongside the existing stream_id. Zero values indicate "no resolved dependency". * CSV: a HIP-specific 10-column encoder (hip_api_csv_encoder) is added so the HIP CSV gains Stream_Id (previously not surfaced), Wait_On_Stream_Id, and Wait_On_Correlation_Id columns appended after the existing 7. Other domains keep the 7-column api_csv_encoder unchanged. Validation ========== A minimal HIP workload (2 streams, hipEventRecord/hipStreamWaitEvent cross- stream edge, plus a separate host-side hipEventSynchronize) was run under rocprofv3 with --hip-trace --output-format json csv. The JSON output has the wait_on_correlation_id field on every hip_api record (default 0) and non-zero on exactly the 2 sync records. The CSV rows for those records show the waiter's stream and the producer's stream + corr_id correctly linked back to the originating hipEventRecord calls. --- .../rocprofiler-sdk/source/lib/output/csv.hpp | 1 + .../source/lib/output/generateCSV.cpp | 20 ++- .../source/lib/output/generatePerfetto.cpp | 86 +++++++--- .../source/lib/output/stream_info.hpp | 27 ++++ .../lib/rocprofiler-sdk-tool/CMakeLists.txt | 5 +- .../event_producer_map.cpp | 141 ++++++++++++++++ .../event_producer_map.hpp | 108 +++++++++++++ .../source/lib/rocprofiler-sdk-tool/tool.cpp | 150 +++++++++++++++++- 8 files changed, 506 insertions(+), 32 deletions(-) create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp diff --git a/projects/rocprofiler-sdk/source/lib/output/csv.hpp b/projects/rocprofiler-sdk/source/lib/output/csv.hpp index 90a7c98f057..02ea83015e9 100644 --- a/projects/rocprofiler-sdk/source/lib/output/csv.hpp +++ b/projects/rocprofiler-sdk/source/lib/output/csv.hpp @@ -100,6 +100,7 @@ struct csv_encoder }; using api_csv_encoder = csv_encoder<7>; +using hip_api_csv_encoder = csv_encoder<10>; using agent_info_csv_encoder = csv_encoder<53>; using counter_collection_csv_encoder = csv_encoder<19>; using memory_allocation_csv_encoder = csv_encoder<8>; diff --git a/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp b/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp index 024da2d44ca..f6933ddf7a2 100644 --- a/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp +++ b/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp @@ -339,23 +339,32 @@ generate_csv(const output_config& cfg, if(cfg.stats && stats) write_stats(get_stats_output_file(cfg, domain_type::HIP), stats.entries); + // HIP CSV is widened to 10 columns: in addition to the standard 7 it carries + // Stream_Id (the call's own stream) and the inter-stream sync metadata + // Wait_On_Stream_Id / Wait_On_Correlation_Id, which are populated only on + // hipStreamWaitEvent / hipEventSynchronize records whose awaited event was + // observed being recorded earlier in the same trace; otherwise the wait_on_* + // columns are 0 (default-stream-id / no-resolved-dependency sentinel). auto ofs = tool::csv_output_file{cfg, domain_type::HIP, - tool::csv::api_csv_encoder{}, + tool::csv::hip_api_csv_encoder{}, {"Domain", "Function", "Process_Id", "Thread_Id", "Correlation_Id", "Start_Timestamp", - "End_Timestamp"}}; + "End_Timestamp", + "Stream_Id", + "Wait_On_Stream_Id", + "Wait_On_Correlation_Id"}}; for(auto ditr : data) { for(auto record : data.get(ditr)) { auto row_ss = std::stringstream{}; auto api_name = tool_metadata.get_operation_name(record.kind, record.operation); - rocprofiler::tool::csv::api_csv_encoder::write_row( + rocprofiler::tool::csv::hip_api_csv_encoder::write_row( row_ss, tool_metadata.get_kind_name(record.kind), api_name, @@ -363,7 +372,10 @@ generate_csv(const output_config& cfg, record.thread_id, record.correlation_id.internal, record.start_timestamp, - record.end_timestamp); + record.end_timestamp, + record.stream_id.handle, + record.wait_on_stream.handle, + record.wait_on_correlation_id); ofs << row_ss.str(); } diff --git a/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp b/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp index b094913f265..a43af89e5e1 100644 --- a/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp +++ b/projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp @@ -366,29 +366,69 @@ write_perfetto( auto name = buffer_names.at(itr.kind, itr.operation); auto& track = thread_tracks.at(itr.thread_id); - TRACE_EVENT_BEGIN(sdk::perfetto_category::name, - ::perfetto::StaticString(name.data()), - 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), - "tid", - itr.thread_id, - "kind", - itr.kind, - "operation", - itr.operation, - "corr_id", - itr.correlation_id.internal, - "ancestor_id", - itr.correlation_id.ancestor, - "stream_ID", - itr.stream_id.handle); + // When this record is a sync API (hipStreamWaitEvent / hipEventSynchronize) + // whose awaited event has a resolved producer, emit a SECOND Flow whose + // id is the producer's correlation_id. Perfetto draws an arrow between + // all events sharing the same flow id, so this connects the waiter slice + // back to the original hipEventRecord call. Two separate macro invocations + // because the variadic flow-arg list has to be fixed at expansion time. + if(itr.wait_on_correlation_id != 0) + { + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(name.data()), + track, + itr.start_timestamp, + ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), + ::perfetto::Flow::ProcessScoped(itr.wait_on_correlation_id), + "begin_ns", + itr.start_timestamp, + "end_ns", + itr.end_timestamp, + "delta_ns", + (itr.end_timestamp - itr.start_timestamp), + "tid", + itr.thread_id, + "kind", + itr.kind, + "operation", + itr.operation, + "corr_id", + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor, + "stream_ID", + itr.stream_id.handle, + "wait_on_stream_ID", + itr.wait_on_stream.handle, + "wait_on_corr_id", + itr.wait_on_correlation_id); + } + else + { + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(name.data()), + 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), + "tid", + itr.thread_id, + "kind", + itr.kind, + "operation", + itr.operation, + "corr_id", + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor, + "stream_ID", + itr.stream_id.handle); + } TRACE_EVENT_END( sdk::perfetto_category::name, track, itr.end_timestamp); diff --git a/projects/rocprofiler-sdk/source/lib/output/stream_info.hpp b/projects/rocprofiler-sdk/source/lib/output/stream_info.hpp index 990c586c0eb..384a8b4eb25 100644 --- a/projects/rocprofiler-sdk/source/lib/output/stream_info.hpp +++ b/projects/rocprofiler-sdk/source/lib/output/stream_info.hpp @@ -120,6 +120,19 @@ struct tool_buffer_tracing_hip_api_ext_record_t : rocprofiler_buffer_tracing_hip , stream_id{_stream_id} {} + // Extended constructor used by the buffer-flush path when the record is for a + // synchronization API (hipStreamWaitEvent / hipEventSynchronize) and the producer + // of the awaited event has been resolved by the HIP_RUNTIME_API callback service. + tool_buffer_tracing_hip_api_ext_record_t(const base_type& _base, + const rocprofiler_stream_id_t _stream_id, + const rocprofiler_stream_id_t _wait_on_stream, + const uint64_t _wait_on_correlation_id) + : base_type{_base} + , stream_id{_stream_id} + , wait_on_stream{_wait_on_stream} + , wait_on_correlation_id{_wait_on_correlation_id} + {} + tool_buffer_tracing_hip_api_ext_record_t() = delete; ~tool_buffer_tracing_hip_api_ext_record_t() = default; tool_buffer_tracing_hip_api_ext_record_t(const tool_buffer_tracing_hip_api_ext_record_t&) = @@ -132,6 +145,18 @@ struct tool_buffer_tracing_hip_api_ext_record_t : rocprofiler_buffer_tracing_hip tool_buffer_tracing_hip_api_ext_record_t&&) noexcept = default; rocprofiler_stream_id_t stream_id = {}; + + // Stream on which the awaited event was last hipEventRecord-ed. + // Default {.handle = 0} (= the default stream) is also used as the sentinel + // for "no resolved dependency" -- distinguished from a real default-stream + // dependency by wait_on_correlation_id != 0 below. + rocprofiler_stream_id_t wait_on_stream = {}; + + // correlation_id of the hipEventRecord call that produced the awaited event. + // 0 means "no resolved dependency" (this record is not a wait-style API, or + // the producer was not observed -- e.g. event was recorded before tracing + // started, or by a different process). + uint64_t wait_on_correlation_id = 0; }; } // namespace tool @@ -174,6 +199,8 @@ save(ArchiveT& ar, const ::rocprofiler::tool::tool_buffer_tracing_hip_api_ext_re { cereal::save(ar, static_cast(data)); SAVE_DATA_FIELD(stream_id); + SAVE_DATA_FIELD(wait_on_stream); + SAVE_DATA_FIELD(wait_on_correlation_id); } #undef SAVE_DATA_FIELD diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt index 9cc12925758..bf2e2291ef5 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt @@ -4,9 +4,10 @@ rocprofiler_activate_clang_tidy() -set(TOOL_HEADERS config.hpp execution_profile.hpp helper.hpp stream_stack.hpp) +set(TOOL_HEADERS config.hpp execution_profile.hpp helper.hpp stream_stack.hpp + event_producer_map.hpp) -set(TOOL_SOURCES config.cpp main.c tool.cpp stream_stack.cpp) +set(TOOL_SOURCES config.cpp main.c tool.cpp stream_stack.cpp event_producer_map.cpp) add_library(rocprofiler-sdk-tool SHARED) target_sources(rocprofiler-sdk-tool PRIVATE ${TOOL_SOURCES} ${TOOL_HEADERS}) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp new file mode 100644 index 00000000000..12c90147878 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp @@ -0,0 +1,141 @@ +// MIT License +// +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#include "event_producer_map.hpp" + +#include "lib/common/static_object.hpp" +#include "lib/common/synchronized.hpp" + +#include + +namespace rocprofiler +{ +namespace tool +{ +namespace event_producer +{ +namespace +{ +// Distinct context types so the two static_object singletons don't collide. +struct producer_map_ctx +{}; +struct resolved_wait_map_ctx +{}; + +using producer_map_t = std::unordered_map; +using resolved_wait_map_t = std::unordered_map; + +using sync_producer_map_t = common::Synchronized; +using sync_resolved_wait_map_t = common::Synchronized; + +sync_producer_map_t& +get_producer_map() +{ + static sync_producer_map_t*& _m = + common::static_object::construct(); + return *_m; +} + +sync_resolved_wait_map_t& +get_resolved_wait_map() +{ + static sync_resolved_wait_map_t*& _m = + common::static_object::construct(); + return *_m; +} +} // namespace + +void +record_event_producer(void* event, + rocprofiler_stream_id_t stream_id, + uint64_t correlation_id) +{ + if(event == nullptr) return; + get_producer_map().wlock([&](producer_map_t& m) { + m[event] = producer_info{stream_id, correlation_id}; + }); +} + +void +forget_event(void* event) +{ + if(event == nullptr) return; + get_producer_map().wlock([&](producer_map_t& m) { m.erase(event); }); +} + +std::optional +lookup_event_producer(void* event) +{ + if(event == nullptr) return std::nullopt; + return get_producer_map().rlock([&](const producer_map_t& m) -> std::optional { + auto it = m.find(event); + if(it == m.end()) return std::nullopt; + return it->second; + }); +} + +void +stash_resolved_wait(uint64_t waiter_correlation_id, + rocprofiler_stream_id_t wait_on_stream, + uint64_t wait_on_correlation_id) +{ + if(waiter_correlation_id == 0) return; + get_resolved_wait_map().wlock([&](resolved_wait_map_t& m) { + m[waiter_correlation_id] = + resolved_wait_info{wait_on_stream, wait_on_correlation_id}; + }); +} + +std::optional +consume_resolved_wait(uint64_t waiter_correlation_id) +{ + if(waiter_correlation_id == 0) return std::nullopt; + return get_resolved_wait_map().wlock( + [&](resolved_wait_map_t& m) -> std::optional { + auto it = m.find(waiter_correlation_id); + if(it == m.end()) return std::nullopt; + auto out = it->second; + m.erase(it); + return out; + }); +} + +counters +get_counters() +{ + counters c{}; + c.producer_map_size = + get_producer_map().rlock([](const producer_map_t& m) { return m.size(); }); + c.resolved_wait_size = + get_resolved_wait_map().rlock([](const resolved_wait_map_t& m) { return m.size(); }); + return c; +} + +void +clear_for_testing() +{ + get_producer_map().wlock([](producer_map_t& m) { m.clear(); }); + get_resolved_wait_map().wlock([](resolved_wait_map_t& m) { m.clear(); }); +} +} // namespace event_producer +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp new file mode 100644 index 00000000000..4e1a11876ff --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp @@ -0,0 +1,108 @@ +// MIT License +// +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include + +#include +#include + +namespace rocprofiler +{ +namespace tool +{ +namespace event_producer +{ +// Per-event producer info recorded at hipEventRecord callback time. +// Looked up at hipStreamWaitEvent / hipEventSynchronize callback time so we +// can resolve which (stream, hipEventRecord-call) the waiter is depending on. +struct producer_info +{ + rocprofiler_stream_id_t stream_id; // stream the event was recorded on + uint64_t correlation_id; // correlation_id of the hipEventRecord call +}; + +// Resolved wait info stashed at hipStreamWaitEvent / hipEventSynchronize +// callback time, keyed by the waiter API call's correlation_id. Drained at +// buffer-flush time when the corresponding tool record is materialized. +struct resolved_wait_info +{ + rocprofiler_stream_id_t wait_on_stream; + uint64_t wait_on_correlation_id; +}; + +// ----- producer map (event -> producer) ----- + +// Record that `event` was last hipEventRecord-ed on `stream_id` by the +// HIP API call with the given `correlation_id`. Overwrites any prior entry +// (an event can be re-recorded multiple times). +void +record_event_producer(void* event, + rocprofiler_stream_id_t stream_id, + uint64_t correlation_id); + +// Remove the producer-map entry for `event`. Called from hipEventDestroy +// callback. Safe to call on an event never recorded (no-op). +void +forget_event(void* event); + +// Look up the producer info for `event`. Returns nullopt if the event has +// never been hipEventRecord-ed (or was already destroyed). +std::optional +lookup_event_producer(void* event); + +// ----- resolved-wait side table (waiter_corr_id -> resolved info) ----- + +// Stash the resolved producer info for a waiter API call (hipStreamWaitEvent +// or hipEventSynchronize), keyed by the waiter's correlation_id. The flush-time +// consumer of this entry is the buffer-record-materialization path in tool.cpp, +// which drains the entry via consume_resolved_wait(). +void +stash_resolved_wait(uint64_t waiter_correlation_id, + rocprofiler_stream_id_t wait_on_stream, + uint64_t wait_on_correlation_id); + +// Consume (look-up + erase) the resolved-wait entry for the given waiter +// API call correlation_id. Returns nullopt if no entry was stashed (the +// waiter resolved to no known producer, or this isn't a wait API record). +std::optional +consume_resolved_wait(uint64_t waiter_correlation_id); + +// ----- diagnostics ----- + +// Counters for tests / diagnostics. Returns the current entry counts; +// not authoritative across racing callbacks but useful for sanity checks. +struct counters +{ + std::size_t producer_map_size; + std::size_t resolved_wait_size; +}; +counters +get_counters(); + +// Clear all state. Test-only. +void +clear_for_testing(); +} // namespace event_producer +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp index 3bae2c1ba05..10b8703077b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -24,6 +24,7 @@ #define _DEFAULT_SOURCE 1 #include "config.hpp" +#include "event_producer_map.hpp" #include "execution_profile.hpp" #include "helper.hpp" #include "stream_stack.hpp" @@ -71,6 +72,8 @@ #include #include #include +#include +#include #include #include #include @@ -792,6 +795,84 @@ hip_stream_display_callback(rocprofiler_callback_tracing_record_t record, common::consume_args(user_data, data); } +// Maintains the cross-stream event dependency state used to populate the +// wait_on_stream / wait_on_correlation_id fields on hipStreamWaitEvent / +// hipEventSynchronize tool records. Registered on a dedicated context that +// filters on the small set of HIP runtime API ops involved in event lifecycle +// and synchronization. Runs only on PHASE_EXIT (after the underlying HIP call +// has completed) so the producer state is consistent with what other threads +// will observe via lookup. +void +hip_event_dependency_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t* user_data, + void* data) +{ + if(record.kind != ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API) return; + if(record.phase != ROCPROFILER_CALLBACK_PHASE_EXIT) return; + + auto* payload = static_cast(record.payload); + if(payload == nullptr) return; + + const auto& args = payload->args; + const auto corr_id = record.correlation_id.internal; + const auto producer = rocprofiler::tool::stream::get_stream_id(); + + switch(record.operation) + { + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord: + tool::event_producer::record_event_producer( + args.hipEventRecord.event, producer, corr_id); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord_spt: + tool::event_producer::record_event_producer( + args.hipEventRecord_spt.event, producer, corr_id); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecordWithFlags: + tool::event_producer::record_event_producer( + args.hipEventRecordWithFlags.event, producer, corr_id); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventDestroy: + tool::event_producer::forget_event(args.hipEventDestroy.event); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize: + { + auto info = + tool::event_producer::lookup_event_producer(args.hipEventSynchronize.event); + if(info) + { + tool::event_producer::stash_resolved_wait( + corr_id, info->stream_id, info->correlation_id); + } + break; + } + case ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent: + { + auto info = + tool::event_producer::lookup_event_producer(args.hipStreamWaitEvent.event); + if(info) + { + tool::event_producer::stash_resolved_wait( + corr_id, info->stream_id, info->correlation_id); + } + break; + } + case ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt: + { + auto info = + tool::event_producer::lookup_event_producer(args.hipStreamWaitEvent_spt.event); + if(info) + { + tool::event_producer::stash_resolved_wait( + corr_id, info->stream_id, info->correlation_id); + } + break; + } + default: break; + } + + common::consume_args(user_data, data); +} + // Stores which runtimes have been initialized in metadata void runtime_initialization_callback(rocprofiler_callback_tracing_record_t record, @@ -1281,9 +1362,30 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, static_cast(header->payload); auto stream_id = get_stream_id(record); - tool::write_ring_buffer( - tool::tool_buffer_tracing_hip_api_ext_record_t{*record, stream_id}, - domain_type::HIP); + + // Drain any cross-stream dependency resolved at HIP_RUNTIME_API + // callback time for this record's correlation id. Only present for + // hipStreamWaitEvent / hipEventSynchronize (and their _spt variants) + // when the awaited event was previously hipEventRecord-ed and the + // producer was observed by our callback service. + auto resolved = + tool::event_producer::consume_resolved_wait(record->correlation_id.internal); + + if(resolved) + { + tool::write_ring_buffer( + tool::tool_buffer_tracing_hip_api_ext_record_t{*record, + stream_id, + resolved->wait_on_stream, + resolved->wait_on_correlation_id}, + domain_type::HIP); + } + else + { + tool::write_ring_buffer( + tool::tool_buffer_tracing_hip_api_ext_record_t{*record, stream_id}, + domain_type::HIP); + } } else if(header->kind == ROCPROFILER_BUFFER_TRACING_RCCL_API) { @@ -2048,6 +2150,7 @@ struct tracing_callbacks_t , cntrl_tracing{cntrl_tracing_callback} , kernel_rename{kernel_rename_callback} , hip_stream{hip_stream_display_callback} + , hip_event_dependency{hip_event_dependency_callback} , callback_tracing{callback_tracing_callback} , buffered_tracing{buffered_tracing_callback} , pc_sampling{pc_sampling_callback} @@ -2063,6 +2166,7 @@ struct tracing_callbacks_t , cntrl_tracing{dummy_callback_tracing_callback} , kernel_rename{dummy_callback_tracing_callback} , hip_stream{dummy_callback_tracing_callback} + , hip_event_dependency{dummy_callback_tracing_callback} , callback_tracing{dummy_callback_tracing_callback} , buffered_tracing{dummy_buffered_tracing_callback} , pc_sampling{dummy_buffered_tracing_callback} @@ -2074,6 +2178,7 @@ struct tracing_callbacks_t const rocprofiler_callback_tracing_cb_t cntrl_tracing = nullptr; const rocprofiler_callback_tracing_cb_t kernel_rename = nullptr; const rocprofiler_callback_tracing_cb_t hip_stream = nullptr; + const rocprofiler_callback_tracing_cb_t hip_event_dependency = nullptr; const rocprofiler_callback_tracing_cb_t callback_tracing = nullptr; const rocprofiler_buffer_tracing_cb_t buffered_tracing = nullptr; const rocprofiler_buffer_tracing_cb_t pc_sampling = nullptr; @@ -2610,6 +2715,45 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) start_context(hip_stream_display_ctx, "hip stream"); + // Track cross-stream event dependencies (hipStreamWaitEvent / hipEventSynchronize) + // so the tool record for waiter API calls can carry the producer stream and the + // hipEventRecord correlation id. Gated on hip_runtime_api_trace because the + // resolved fields are only consumed at buffer-flush time for HIP API records, + // and we want zero overhead when HIP tracing is off. + if(tool::get_config().hip_runtime_api_trace && + tool::get_config().benchmark_mode != tool::config::benchmark::sdk_callback_overhead) + { + auto hip_event_dep_ctx = rocprofiler_context_id_t{0}; + + ROCPROFILER_CALL(rocprofiler_create_context(&hip_event_dep_ctx), + "failed to create hip event dependency context"); + + // Filter on just the operations involved in event lifecycle and synchronization. + // Both default and _spt variants are listed because each is dispatched separately + // by the SDK and the args layout differs slightly. + static const auto hip_event_dep_ops = + std::array{ + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord, + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord_spt, + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecordWithFlags, + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventDestroy, + ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize, + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent, + ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt}; + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service( + hip_event_dep_ctx, + ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API, + hip_event_dep_ops.data(), + hip_event_dep_ops.size(), + callbacks.hip_event_dependency, + nullptr), + "hip event dependency tracing configure failed"); + + start_context(hip_event_dep_ctx, "hip event dependency"); + } + // Track if HIP runtime has been initialized via runtime_intialization service auto runtime_initialization_ctx = rocprofiler_context_id_t{0}; From 326362e69acd5964ce996f17901761b4054a6b9a Mon Sep 17 00:00:00 2001 From: ajassani Date: Sun, 24 May 2026 13:32:17 +0000 Subject: [PATCH 2/3] rocprofiler-sdk-tool: tests for inter-stream event sync metadata Adds unit and integration coverage for the wait_on_stream / wait_on_correlation_id fields introduced in the previous commit. Unit tests (source/lib/tests/tool/event_producer_map.cpp, 13 cases): - record / lookup round-trip, including overwrite semantics - null/zero sentinel handling (no-op, never inserted) - forget removes producer entries; lookup of unknown returns nullopt - stash + consume drains the resolved-wait side-table (single-shot) - clear_for_testing wipes both maps - 4-writer x 1000-op stress to surface any locking regression The test binary compiles event_producer_map.cpp directly into the gtest executable (no link to librocprofiler-sdk-tool.so) so each test owns its own translation-unit-local static maps and the suite stays hermetic. Wired into source/lib/tests/CMakeLists.txt via add_subdirectory(tool). Integration test (tests/rocprofv3/hip-interstream-sync, 5 pytest cases driven by rocprofv3 + a HIP test app): The workload (tests/bin/hip-interstream-sync/main.cpp) issues a cross-stream dependency streamA: kernel -> hipEventRecord(eAB) -> kernel | v hipStreamWaitEvent(streamB, eAB) streamB: wait -> kernel and a separate hipEventRecord/hipEventSynchronize pair, then destroys all events and streams. The validator runs rocprofv3 --hip-trace --kernel-trace with json, csv, and pftrace outputs against that binary, then asserts: - the HIP CSV exposes the new Stream_Id, Wait_On_Stream_Id, and Wait_On_Correlation_Id columns - every hipStreamWaitEvent and hipEventSynchronize row carries a non-zero Wait_On_Correlation_Id whose value matches some prior hipEventRecord row's Correlation_Id (no dangling references) - waiter rows report a Wait_On_Stream_Id distinct from their own Stream_Id (cross-stream is the entire point) - non-wait/sync rows still carry the 0 sentinel (no leakage) - the JSON hip_api records carry wait_on_stream and wait_on_correlation_id with the same cross-checks - resolved-dependency counts agree between the CSV and JSON output formats for the same trace Both test apps and the validator are wired into tests/bin/CMakeLists.txt and tests/rocprofv3/CMakeLists.txt. --- .../source/lib/tests/CMakeLists.txt | 1 + .../source/lib/tests/tool/CMakeLists.txt | 39 ++++ .../lib/tests/tool/event_producer_map.cpp | 215 ++++++++++++++++++ .../rocprofiler-sdk/tests/bin/CMakeLists.txt | 1 + .../bin/hip-interstream-sync/CMakeLists.txt | 39 ++++ .../tests/bin/hip-interstream-sync/main.cpp | 117 ++++++++++ .../tests/rocprofv3/CMakeLists.txt | 1 + .../hip-interstream-sync/CMakeLists.txt | 41 ++++ .../hip-interstream-sync/conftest.py | 58 +++++ .../rocprofv3/hip-interstream-sync/pytest.ini | 5 + .../hip-interstream-sync/validate.py | 201 ++++++++++++++++ 11 files changed, 718 insertions(+) create mode 100644 projects/rocprofiler-sdk/source/lib/tests/tool/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp create mode 100644 projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/main.cpp create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/conftest.py create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/validate.py diff --git a/projects/rocprofiler-sdk/source/lib/tests/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/tests/CMakeLists.txt index d170f7a5ece..af0fabe2d9c 100644 --- a/projects/rocprofiler-sdk/source/lib/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/tests/CMakeLists.txt @@ -4,3 +4,4 @@ add_subdirectory(buffering) add_subdirectory(common) add_subdirectory(codeobj) +add_subdirectory(tool) diff --git a/projects/rocprofiler-sdk/source/lib/tests/tool/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/tests/tool/CMakeLists.txt new file mode 100644 index 00000000000..532b159f1b2 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/tests/tool/CMakeLists.txt @@ -0,0 +1,39 @@ +# +# Tests for the rocprofiler-sdk-tool library. +# Implementation files from source/lib/rocprofiler-sdk-tool/ are pulled in +# directly (rather than linking the LD_PRELOAD-shaped librocprofiler-sdk-tool.so) +# so the tests get their own translation-unit-local static state and stay +# hermetic. +# +project(rocprofiler-sdk-tests-tool LANGUAGES C CXX) + +set(tool_test_sources event_producer_map.cpp) + +# Resolve the tool implementation dir relative to this CMakeLists. We avoid +# CMAKE_SOURCE_DIR / PROJECT_SOURCE_DIR here because either may be redefined +# depending on which superproject embeds rocprofiler-sdk. +get_filename_component(_tool_impl_dir + "${CMAKE_CURRENT_LIST_DIR}/../../rocprofiler-sdk-tool" ABSOLUTE) + +set(tool_impl_sources ${_tool_impl_dir}/event_producer_map.cpp) + +add_executable(tool-tests) +target_sources(tool-tests PRIVATE ${tool_test_sources} ${tool_impl_sources}) + +target_include_directories(tool-tests PRIVATE ${_tool_impl_dir}) + +target_link_libraries( + tool-tests + PRIVATE rocprofiler-sdk::rocprofiler-sdk-headers + rocprofiler-sdk::rocprofiler-sdk-common-library + rocprofiler-sdk::rocprofiler-sdk-build-flags + GTest::gtest + GTest::gtest_main) + +rocprofiler_add_unit_test( + TARGET tool-tests + SOURCES ${tool_test_sources} + ENVIRONMENT + "TEST_LOG_LEVEL=info" + "LD_LIBRARY_PATH=${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}:${CMAKE_INSTALL_PREFIX}/lib:${CMAKE_INSTALL_PREFIX}/llvm/lib:${ROCM_PATH}/lib:${ROCM_PATH}/llvm/lib:$ENV{LD_LIBRARY_PATH}" + ) diff --git a/projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp b/projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp new file mode 100644 index 00000000000..e1cf4564cec --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp @@ -0,0 +1,215 @@ +// MIT License +// +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "lib/rocprofiler-sdk-tool/event_producer_map.hpp" + +#include "lib/common/logging.hpp" + +#include + +#include +#include +#include + +namespace +{ +bool _event_producer_map_test_init_logging = + (rocprofiler::common::init_logging("TEST"), true); + +namespace ep = rocprofiler::tool::event_producer; + +void* +fake_event(uintptr_t v) +{ + // Helper to fabricate a non-null hipEvent_t-shaped pointer without ever + // dereferencing it. The map only uses the pointer value as a key. + return reinterpret_cast(v); +} + +rocprofiler_stream_id_t +sid(uint64_t v) +{ + return rocprofiler_stream_id_t{v}; +} + +class EventProducerMap : public ::testing::Test +{ +protected: + void SetUp() override { ep::clear_for_testing(); } + void TearDown() override { ep::clear_for_testing(); } +}; +} // namespace + +TEST_F(EventProducerMap, lookup_unknown_event_returns_nullopt) +{ + EXPECT_FALSE(ep::lookup_event_producer(fake_event(0xdead)).has_value()); +} + +TEST_F(EventProducerMap, lookup_null_event_returns_nullopt) +{ + EXPECT_FALSE(ep::lookup_event_producer(nullptr).has_value()); +} + +TEST_F(EventProducerMap, record_then_lookup_round_trips) +{ + auto* e = fake_event(0x100); + ep::record_event_producer(e, sid(7), 42); + + auto info = ep::lookup_event_producer(e); + ASSERT_TRUE(info.has_value()); + EXPECT_EQ(info->stream_id.handle, 7u); + EXPECT_EQ(info->correlation_id, 42u); +} + +TEST_F(EventProducerMap, record_overwrites_prior_entry) +{ + // Same event re-recorded with newer stream/corr; lookup sees the latest. + auto* e = fake_event(0x200); + ep::record_event_producer(e, sid(1), 10); + ep::record_event_producer(e, sid(2), 20); + + auto info = ep::lookup_event_producer(e); + ASSERT_TRUE(info.has_value()); + EXPECT_EQ(info->stream_id.handle, 2u); + EXPECT_EQ(info->correlation_id, 20u); +} + +TEST_F(EventProducerMap, record_null_event_is_silent_noop) +{ + ep::record_event_producer(nullptr, sid(1), 5); + EXPECT_EQ(ep::get_counters().producer_map_size, 0u); +} + +TEST_F(EventProducerMap, forget_event_removes_entry) +{ + auto* e = fake_event(0x300); + ep::record_event_producer(e, sid(3), 11); + ASSERT_TRUE(ep::lookup_event_producer(e).has_value()); + + ep::forget_event(e); + EXPECT_FALSE(ep::lookup_event_producer(e).has_value()); +} + +TEST_F(EventProducerMap, forget_unknown_event_is_silent_noop) +{ + ep::forget_event(fake_event(0x400)); + EXPECT_EQ(ep::get_counters().producer_map_size, 0u); +} + +TEST_F(EventProducerMap, distinct_events_are_independent) +{ + auto* e1 = fake_event(0x500); + auto* e2 = fake_event(0x501); + ep::record_event_producer(e1, sid(1), 100); + ep::record_event_producer(e2, sid(2), 200); + + auto i1 = ep::lookup_event_producer(e1); + auto i2 = ep::lookup_event_producer(e2); + + ASSERT_TRUE(i1.has_value()); + ASSERT_TRUE(i2.has_value()); + EXPECT_EQ(i1->stream_id.handle, 1u); + EXPECT_EQ(i1->correlation_id, 100u); + EXPECT_EQ(i2->stream_id.handle, 2u); + EXPECT_EQ(i2->correlation_id, 200u); + + ep::forget_event(e1); + EXPECT_FALSE(ep::lookup_event_producer(e1).has_value()); + EXPECT_TRUE(ep::lookup_event_producer(e2).has_value()); +} + +TEST_F(EventProducerMap, stash_then_consume_resolved_wait_drains) +{ + ep::stash_resolved_wait(/*waiter_corr_id=*/55, sid(9), /*producer_corr_id=*/22); + + auto r = ep::consume_resolved_wait(55); + ASSERT_TRUE(r.has_value()); + EXPECT_EQ(r->wait_on_stream.handle, 9u); + EXPECT_EQ(r->wait_on_correlation_id, 22u); + + // Second consume on the same key returns nullopt: stash is single-shot. + EXPECT_FALSE(ep::consume_resolved_wait(55).has_value()); +} + +TEST_F(EventProducerMap, consume_unknown_corr_id_returns_nullopt) +{ + EXPECT_FALSE(ep::consume_resolved_wait(999).has_value()); +} + +TEST_F(EventProducerMap, stash_zero_corr_id_is_silent_noop) +{ + // 0 is reserved as the "no-resolved-dependency" sentinel. + ep::stash_resolved_wait(0, sid(1), 1); + EXPECT_EQ(ep::get_counters().resolved_wait_size, 0u); + EXPECT_FALSE(ep::consume_resolved_wait(0).has_value()); +} + +TEST_F(EventProducerMap, clear_for_testing_resets_both_maps) +{ + ep::record_event_producer(fake_event(0x600), sid(1), 1); + ep::stash_resolved_wait(7, sid(2), 2); + auto pre = ep::get_counters(); + EXPECT_EQ(pre.producer_map_size, 1u); + EXPECT_EQ(pre.resolved_wait_size, 1u); + + ep::clear_for_testing(); + + auto post = ep::get_counters(); + EXPECT_EQ(post.producer_map_size, 0u); + EXPECT_EQ(post.resolved_wait_size, 0u); +} + +TEST_F(EventProducerMap, concurrent_record_and_forget_does_not_corrupt) +{ + // Stress test the shared/unique lock pairing under contention. + // We don't assert specific outcomes for racy operations -- only that the + // map remains internally consistent (no crash, lookups return either a + // valid entry or nullopt cleanly). + constexpr std::size_t kWriters = 4; + constexpr std::size_t kPerThread = 1000; + + std::atomic start{false}; + std::vector writers; + writers.reserve(kWriters); + for(std::size_t t = 0; t < kWriters; ++t) + { + writers.emplace_back([t, &start]() { + while(!start.load(std::memory_order_acquire)) + { + std::this_thread::yield(); + } + for(std::size_t i = 0; i < kPerThread; ++i) + { + auto* e = fake_event(0x10000 + (t * kPerThread) + i); + ep::record_event_producer(e, sid(t + 1), i); + auto r = ep::lookup_event_producer(e); + EXPECT_TRUE(r.has_value()); + ep::forget_event(e); + } + }); + } + start.store(true, std::memory_order_release); + for(auto& th : writers) th.join(); + + // After all writers complete, all keys they touched should be erased. + EXPECT_EQ(ep::get_counters().producer_map_size, 0u); +} diff --git a/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt index e50f3ae645c..2886654a98d 100644 --- a/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt @@ -48,3 +48,4 @@ add_subdirectory(module-loading-test) add_subdirectory(late-start-tracing) add_subdirectory(thread-trace) add_subdirectory(hip-graph-bubbles) +add_subdirectory(hip-interstream-sync) diff --git a/projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/CMakeLists.txt new file mode 100644 index 00000000000..f38b6b9d4b9 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/CMakeLists.txt @@ -0,0 +1,39 @@ +# +# Test app for inter-stream HIP event synchronization metadata. +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +if(NOT CMAKE_HIP_COMPILER) + find_program( + amdclangpp_EXECUTABLE + NAMES amdclang++ + HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATH_SUFFIXES bin llvm/bin NO_CACHE) + mark_as_advanced(amdclangpp_EXECUTABLE) + + if(amdclangpp_EXECUTABLE) + set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}") + endif() +endif() + +project(rocprofiler-sdk-tests-bin-hip-interstream-sync LANGUAGES CXX HIP) + +foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO) + if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "") + set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}") + endif() +endforeach() + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_HIP_STANDARD 17) +set(CMAKE_HIP_EXTENSIONS OFF) +set(CMAKE_HIP_STANDARD_REQUIRED ON) + +set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP) +add_executable(hip-interstream-sync) +target_sources(hip-interstream-sync PRIVATE main.cpp) +target_link_libraries(hip-interstream-sync + PRIVATE rocprofiler-sdk::tests-build-flags) diff --git a/projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/main.cpp b/projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/main.cpp new file mode 100644 index 00000000000..adb188f8016 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/bin/hip-interstream-sync/main.cpp @@ -0,0 +1,117 @@ +// MIT License +// +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +// Minimal cross-stream sync workload that exercises rocprofv3's new +// wait_on_stream / wait_on_correlation_id metadata on hipStreamWaitEvent +// and hipEventSynchronize records. +// +// Pattern: +// stream A: [ kernelA ] -- hipEventRecord(e) -- [ kernelA2 ] +// | +// v hipStreamWaitEvent(stream B, e) +// stream B: [ wait ] -- [ kernelB ] +// +// Then a separate hipEventSynchronize on the host exercises that path, and +// hipEventDestroy at the end causes the producer map to forget the entry. + +#include + +#include +#include + +#define HIP_CHECK(expr) \ + do \ + { \ + hipError_t _e = (expr); \ + if(_e != hipSuccess) \ + { \ + std::fprintf(stderr, \ + "%s:%d %s: %s\n", \ + __FILE__, \ + __LINE__, \ + #expr, \ + hipGetErrorString(_e)); \ + std::exit(1); \ + } \ + } while(0) + +__global__ void +busy_kernel(float* buf, int n, int iters) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if(i < n) + { + float v = buf[i]; + for(int k = 0; k < iters; ++k) + v = v * 1.000001f + 0.0001f; + buf[i] = v; + } +} + +int +main() +{ + constexpr int N = 1 << 16; + constexpr int ITERS = 5000; + + float* d_a = nullptr; + float* d_b = nullptr; + HIP_CHECK(hipMalloc(&d_a, N * sizeof(float))); + HIP_CHECK(hipMalloc(&d_b, N * sizeof(float))); + HIP_CHECK(hipMemset(d_a, 0, N * sizeof(float))); + HIP_CHECK(hipMemset(d_b, 0, N * sizeof(float))); + + hipStream_t sA = nullptr; + hipStream_t sB = nullptr; + HIP_CHECK(hipStreamCreate(&sA)); + HIP_CHECK(hipStreamCreate(&sB)); + + hipEvent_t eAB = nullptr; + hipEvent_t eHost = nullptr; + HIP_CHECK(hipEventCreate(&eAB)); + HIP_CHECK(hipEventCreate(&eHost)); + + dim3 grid((N + 255) / 256), block(256); + + // cross-stream dependency: A -> wait -> B + busy_kernel<<>>(d_a, N, ITERS); + HIP_CHECK(hipEventRecord(eAB, sA)); + HIP_CHECK(hipStreamWaitEvent(sB, eAB, 0)); + busy_kernel<<>>(d_b, N, ITERS); + + // host-side sync on a separate event + busy_kernel<<>>(d_a, N, ITERS); + HIP_CHECK(hipEventRecord(eHost, sA)); + HIP_CHECK(hipEventSynchronize(eHost)); + + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipEventDestroy(eAB)); + HIP_CHECK(hipEventDestroy(eHost)); + HIP_CHECK(hipStreamDestroy(sA)); + HIP_CHECK(hipStreamDestroy(sB)); + HIP_CHECK(hipFree(d_a)); + HIP_CHECK(hipFree(d_b)); + + std::printf("hip-interstream-sync OK\n"); + return 0; +} diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt index 17a61673b78..8296432b67e 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt @@ -56,3 +56,4 @@ add_subdirectory(kfd) add_subdirectory(roctx-pause-resume) add_subdirectory(mpi-ranks) add_subdirectory(hip-graph-bubbles-test) +add_subdirectory(hip-interstream-sync) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/CMakeLists.txt new file mode 100644 index 00000000000..5de3223d67d --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/CMakeLists.txt @@ -0,0 +1,41 @@ +# +# rocprofv3 integration test: HIP inter-stream event-sync metadata. +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-sdk-tests-rocprofv3-hip-interstream-sync + LANGUAGES CXX + VERSION 0.0.0) + +string(REPLACE "LD_PRELOAD=" "--preload;" PRELOAD_ARGS + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") + +find_package(rocprofiler-sdk REQUIRED) + +# Run the cross-stream workload under rocprofv3 with HIP runtime tracing +# and dump JSON + CSV + pftrace so the validator can inspect all three. +rocprofiler_add_integration_execute_test( + rocprofv3-test-hip-interstream-sync + COMMAND + $ --hip-trace --kernel-trace -d + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out --output-format json csv pftrace + ${PRELOAD_ARGS} -- $ + DEPENDS hip-interstream-sync + TIMEOUT 60 + LABELS "integration-tests" + FIXTURES_SETUP rocprofv3-test-hip-interstream-sync + FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") + +rocprofiler_add_integration_validate_test( + rocprofv3-test-hip-interstream-sync + TEST_PATHS validate.py + COPY conftest.py + CONFIG pytest.ini + ARGS --hip-trace-input + ${CMAKE_CURRENT_BINARY_DIR}/hip-interstream-sync-trace/out_hip_api_trace.csv + --json-input + ${CMAKE_CURRENT_BINARY_DIR}/hip-interstream-sync-trace/out_results.json + TIMEOUT 45 + LABELS "integration-tests" + FIXTURES_REQUIRED rocprofv3-test-hip-interstream-sync) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/conftest.py new file mode 100644 index 00000000000..85784c4fc40 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/conftest.py @@ -0,0 +1,58 @@ +#!/usr/bin/env python3 + +# MIT License +# +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import csv +import json + +import pytest + +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list + + +def pytest_addoption(parser): + parser.addoption( + "--hip-trace-input", + action="store", + help="Path to HIP API tracing CSV file.", + ) + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) + + +@pytest.fixture +def hip_trace_csv_rows(request): + filename = request.config.getoption("--hip-trace-input") + with open(filename, "r") as inp: + return list(csv.DictReader(inp)) + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/pytest.ini new file mode 100644 index 00000000000..8bf72b0989b --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/validate.py new file mode 100644 index 00000000000..e0779afd020 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hip-interstream-sync/validate.py @@ -0,0 +1,201 @@ +#!/usr/bin/env python3 + +# MIT License +# +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +# End-to-end check that rocprofv3 emits inter-stream-synchronization +# metadata (wait_on_stream + wait_on_correlation_id) on HIP runtime API +# records for hipStreamWaitEvent and hipEventSynchronize, and that the +# referenced producer correlation IDs actually point at hipEventRecord +# rows earlier in the same trace. + +import sys + +import pytest + + +# ----------------------------------------------------------------------------- +# CSV checks +# ----------------------------------------------------------------------------- + + +def _hip_csv_columns_present(rows): + # If the trace is empty we have a different problem, surface that first. + assert rows, "HIP API CSV had zero rows" + cols = set(rows[0].keys()) + # The new tri-tuple is what this feature adds. + for required in ("Stream_Id", "Wait_On_Stream_Id", "Wait_On_Correlation_Id"): + assert required in cols, f"missing CSV column: {required} (have {sorted(cols)})" + + +def test_hip_csv_has_inter_stream_columns(hip_trace_csv_rows): + _hip_csv_columns_present(hip_trace_csv_rows) + + +def test_hip_csv_wait_dependencies_resolve(hip_trace_csv_rows): + _hip_csv_columns_present(hip_trace_csv_rows) + + producer_corr_ids = { + int(r["Correlation_Id"]) for r in hip_trace_csv_rows + if r["Function"] == "hipEventRecord" + } + assert producer_corr_ids, "workload was supposed to call hipEventRecord at least once" + + waiters = [r for r in hip_trace_csv_rows if r["Function"] == "hipStreamWaitEvent"] + syncs = [r for r in hip_trace_csv_rows if r["Function"] == "hipEventSynchronize"] + assert waiters, "workload was supposed to call hipStreamWaitEvent at least once" + assert syncs, "workload was supposed to call hipEventSynchronize at least once" + + # Every hipStreamWaitEvent we emitted in the workload had a matching prior + # hipEventRecord, so all waiter rows should carry a resolved dependency. + resolved_waiters = [w for w in waiters if int(w["Wait_On_Correlation_Id"]) != 0] + assert resolved_waiters, ( + "no hipStreamWaitEvent row had a non-zero Wait_On_Correlation_Id; " + "the producer-map lookup did not resolve any edges" + ) + + for w in resolved_waiters: + wait_corr = int(w["Wait_On_Correlation_Id"]) + assert wait_corr in producer_corr_ids, ( + f"Wait_On_Correlation_Id {wait_corr} does not match any " + f"hipEventRecord Correlation_Id; resolution pointed at a " + f"non-producer call" + ) + # Cross-stream wait: the waiter's own stream must differ from the + # producer's stream, otherwise the dependency tells us nothing. + assert int(w["Wait_On_Stream_Id"]) != int(w["Stream_Id"]), ( + f"hipStreamWaitEvent reports waiting on its own stream " + f"(stream={w['Stream_Id']}); the workload always waits cross-stream" + ) + + resolved_syncs = [s for s in syncs if int(s["Wait_On_Correlation_Id"]) != 0] + assert resolved_syncs, ( + "no hipEventSynchronize row had a non-zero Wait_On_Correlation_Id" + ) + for s in resolved_syncs: + wait_corr = int(s["Wait_On_Correlation_Id"]) + assert wait_corr in producer_corr_ids, ( + f"hipEventSynchronize Wait_On_Correlation_Id {wait_corr} does not " + f"match any hipEventRecord Correlation_Id" + ) + + # Non-dependency rows must still carry the sentinel 0, never garbage. + for r in hip_trace_csv_rows: + if r["Function"] not in ("hipStreamWaitEvent", "hipEventSynchronize"): + assert int(r["Wait_On_Correlation_Id"]) == 0, ( + f"row {r['Function']} carries Wait_On_Correlation_Id=" + f"{r['Wait_On_Correlation_Id']}; only wait/sync APIs should " + f"ever resolve a dependency" + ) + + +# ----------------------------------------------------------------------------- +# JSON checks +# ----------------------------------------------------------------------------- + + +def _iter_hip_api_records(json_data): + data = json_data["rocprofiler-sdk-tool"] + return data["buffer_records"]["hip_api"] + + +def _op_name(json_data, kind, op): + data = json_data["rocprofiler-sdk-tool"] + return data["strings"]["buffer_records"][kind]["operations"][op] + + +def test_hip_json_has_inter_stream_fields(json_data): + records = _iter_hip_api_records(json_data) + assert records, "no hip_api records in JSON output" + sample = records[0] + assert "wait_on_stream" in sample, ( + f"hip_api record missing wait_on_stream field; " + f"keys present: {sorted(sample.keys())}" + ) + assert "wait_on_correlation_id" in sample, ( + f"hip_api record missing wait_on_correlation_id field; " + f"keys present: {sorted(sample.keys())}" + ) + # wait_on_stream serializes as the nested rocprofiler_stream_id_t cereal + # struct, so it has a "handle" subfield. + assert "handle" in sample["wait_on_stream"], ( + f"wait_on_stream should serialize as {{'handle': N}}, got " + f"{sample['wait_on_stream']!r}" + ) + + +def test_hip_json_wait_dependencies_resolve(json_data): + records = _iter_hip_api_records(json_data) + + producer_corr_ids = set() + waiters = [] + syncs = [] + for rec in records: + name = _op_name(json_data, rec["kind"], rec["operation"]) + if name == "hipEventRecord": + producer_corr_ids.add(rec["correlation_id"]["internal"]) + elif name == "hipStreamWaitEvent": + waiters.append(rec) + elif name == "hipEventSynchronize": + syncs.append(rec) + + assert producer_corr_ids, "expected at least one hipEventRecord in JSON" + assert waiters, "expected at least one hipStreamWaitEvent in JSON" + assert syncs, "expected at least one hipEventSynchronize in JSON" + + resolved_waiters = [w for w in waiters if w["wait_on_correlation_id"] != 0] + assert resolved_waiters, "no JSON hipStreamWaitEvent resolved a dependency" + for w in resolved_waiters: + assert w["wait_on_correlation_id"] in producer_corr_ids + assert w["wait_on_stream"]["handle"] != w["stream_id"]["handle"] + + resolved_syncs = [s for s in syncs if s["wait_on_correlation_id"] != 0] + assert resolved_syncs, "no JSON hipEventSynchronize resolved a dependency" + for s in resolved_syncs: + assert s["wait_on_correlation_id"] in producer_corr_ids + + +def test_csv_and_json_agree_on_resolved_count(hip_trace_csv_rows, json_data): + # Same workload, same tool run, two output formats: counts of resolved + # dependencies on the two relevant API calls must match exactly. + csv_resolved = sum( + 1 + for r in hip_trace_csv_rows + if r["Function"] in ("hipStreamWaitEvent", "hipEventSynchronize") + and int(r["Wait_On_Correlation_Id"]) != 0 + ) + json_resolved = sum( + 1 + for rec in _iter_hip_api_records(json_data) + if _op_name(json_data, rec["kind"], rec["operation"]) + in ("hipStreamWaitEvent", "hipEventSynchronize") + and rec["wait_on_correlation_id"] != 0 + ) + assert csv_resolved == json_resolved, ( + f"resolved-dependency count disagrees between CSV ({csv_resolved}) " + f"and JSON ({json_resolved}) for the same trace" + ) + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) From 44b72a0a469ac816916ac3f9a2a40d2365d365b3 Mon Sep 17 00:00:00 2001 From: ajassani Date: Mon, 25 May 2026 22:26:05 +0000 Subject: [PATCH 3/3] rocprofiler-sdk-tool: tighten inter-stream sync dependency tracking Pre-review polish on the inter-stream event-sync metadata feature so the upstream review focuses on the design rather than mechanical concerns. Hot-path guard on the buffer-flush side-table lookup ---------------------------------------------------- buffered_tracing_callback was unconditionally calling event_producer::consume_resolved_wait() for every HIP API buffer record. The side-table is only ever populated for hipStreamWaitEvent / hipStreamWaitEvent_spt / hipEventSynchronize records, so the lookup was a guaranteed miss (with a write-lock acquisition) on ~99.95% of records on a representative DDP workload. Gate the lookup on the record's operation kind so the cost on every other HIP API call is zero. Exception safety in the callback -------------------------------- The dependency callback runs synchronously inside the SDK dispatch path and mutates std::unordered_map state, which can throw std::bad_alloc under memory pressure. An uncaught exception would propagate into the user's HIP call site. Wrap the callback body in a try/catch and log via ROCP_WARNING on the unhappy path. Reattach hygiene + log visibility --------------------------------- The producer / resolved-wait maps live in common::static_object and persist across tool_init -> tool_detach -> tool_init cycles. hipEvent_t pointers can be reused across sessions, which could cause false-positive dependency resolutions on reattach. Clear both maps at tool_init right before the dependency context is created. Also emit a one-shot ROCP_INFO so users can see the service was enabled in the tool init log. Header tidy ----------- Rename event_producer::clear_for_testing() to clear_all() and document that it is also used at tool_init for reattach state. The old name was misleading now that the function has a production caller. Validation ---------- * tool-tests (gtest): 13/13 pass, including the renamed EventProducerMap.clear_all_resets_both_maps case. * rocprofv3-test-hip-interstream-sync (integration): 6/6 pass (1 execute + 5 validate). * Re-ran the 2-rank DDP workload under rocprofv3 with --hip-trace --kernel-trace --output-format json csv -- same resolved-dependency counts as before the guard (149,872 HIP API records, 246 with a non-zero Stream_Id, 78 with a non-zero Wait_On_Correlation_Id); metadata fidelity unchanged, just less locking on the hot path. --- .../event_producer_map.cpp | 2 +- .../event_producer_map.hpp | 6 +- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 136 +++++++++++------- .../lib/tests/tool/event_producer_map.cpp | 8 +- 4 files changed, 90 insertions(+), 62 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp index 12c90147878..6d76ae263ad 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.cpp @@ -131,7 +131,7 @@ get_counters() } void -clear_for_testing() +clear_all() { get_producer_map().wlock([](producer_map_t& m) { m.clear(); }); get_resolved_wait_map().wlock([](resolved_wait_map_t& m) { m.clear(); }); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp index 4e1a11876ff..58c65e9b30a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/event_producer_map.hpp @@ -100,9 +100,11 @@ struct counters counters get_counters(); -// Clear all state. Test-only. +// Clear all producer + resolved-wait state. Called from tool_init so a +// reattach cycle starts with an empty map (event pointers can be reused +// across sessions), and from unit tests to reset state between cases. void -clear_for_testing(); +clear_all(); } // namespace event_producer } // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp index 10b8703077b..2e4e512468e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -813,61 +813,72 @@ hip_event_dependency_callback(rocprofiler_callback_tracing_record_t record, auto* payload = static_cast(record.payload); if(payload == nullptr) return; - const auto& args = payload->args; - const auto corr_id = record.correlation_id.internal; - const auto producer = rocprofiler::tool::stream::get_stream_id(); - - switch(record.operation) - { - case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord: - tool::event_producer::record_event_producer( - args.hipEventRecord.event, producer, corr_id); - break; - case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord_spt: - tool::event_producer::record_event_producer( - args.hipEventRecord_spt.event, producer, corr_id); - break; - case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecordWithFlags: - tool::event_producer::record_event_producer( - args.hipEventRecordWithFlags.event, producer, corr_id); - break; - case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventDestroy: - tool::event_producer::forget_event(args.hipEventDestroy.event); - break; - case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize: - { - auto info = - tool::event_producer::lookup_event_producer(args.hipEventSynchronize.event); - if(info) + // Wrap the side-table updates in a top-level catch. The underlying + // unordered_map ops can throw std::bad_alloc under memory pressure and + // this callback runs synchronously inside the SDK dispatch path -- an + // uncaught exception would propagate into the user's HIP call site. + try + { + const auto& args = payload->args; + const auto corr_id = record.correlation_id.internal; + const auto producer = rocprofiler::tool::stream::get_stream_id(); + + switch(record.operation) + { + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord: + tool::event_producer::record_event_producer( + args.hipEventRecord.event, producer, corr_id); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecord_spt: + tool::event_producer::record_event_producer( + args.hipEventRecord_spt.event, producer, corr_id); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventRecordWithFlags: + tool::event_producer::record_event_producer( + args.hipEventRecordWithFlags.event, producer, corr_id); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventDestroy: + tool::event_producer::forget_event(args.hipEventDestroy.event); + break; + case ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize: { - tool::event_producer::stash_resolved_wait( - corr_id, info->stream_id, info->correlation_id); + auto info = + tool::event_producer::lookup_event_producer(args.hipEventSynchronize.event); + if(info) + { + tool::event_producer::stash_resolved_wait( + corr_id, info->stream_id, info->correlation_id); + } + break; } - break; - } - case ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent: - { - auto info = - tool::event_producer::lookup_event_producer(args.hipStreamWaitEvent.event); - if(info) + case ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent: { - tool::event_producer::stash_resolved_wait( - corr_id, info->stream_id, info->correlation_id); + auto info = + tool::event_producer::lookup_event_producer(args.hipStreamWaitEvent.event); + if(info) + { + tool::event_producer::stash_resolved_wait( + corr_id, info->stream_id, info->correlation_id); + } + break; } - break; - } - case ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt: - { - auto info = - tool::event_producer::lookup_event_producer(args.hipStreamWaitEvent_spt.event); - if(info) + case ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt: { - tool::event_producer::stash_resolved_wait( - corr_id, info->stream_id, info->correlation_id); + auto info = + tool::event_producer::lookup_event_producer(args.hipStreamWaitEvent_spt.event); + if(info) + { + tool::event_producer::stash_resolved_wait( + corr_id, info->stream_id, info->correlation_id); + } + break; } - break; + default: break; } - default: break; + } catch(const std::exception& e) { + ROCP_WARNING << "hip_event_dependency_callback dropped a record: " << e.what(); + } catch(...) { + ROCP_WARNING << "hip_event_dependency_callback dropped a record: unknown exception"; } common::consume_args(user_data, data); @@ -1363,13 +1374,20 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, auto stream_id = get_stream_id(record); - // Drain any cross-stream dependency resolved at HIP_RUNTIME_API - // callback time for this record's correlation id. Only present for - // hipStreamWaitEvent / hipEventSynchronize (and their _spt variants) - // when the awaited event was previously hipEventRecord-ed and the - // producer was observed by our callback service. - auto resolved = - tool::event_producer::consume_resolved_wait(record->correlation_id.internal); + // Hot path: this branch runs for every HIP API buffer record. + // Only the three sync APIs can ever carry a resolved dependency + // in the side-table, so gate the (locking) map lookup on the + // record's operation kind. On a representative DDP workload this + // skips the lookup for ~99.95% of records. + const auto op = record->operation; + std::optional resolved; + if(op == ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent || + op == ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamWaitEvent_spt || + op == ROCPROFILER_HIP_RUNTIME_API_ID_hipEventSynchronize) + { + resolved = tool::event_producer::consume_resolved_wait( + record->correlation_id.internal); + } if(resolved) { @@ -2723,6 +2741,14 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) if(tool::get_config().hip_runtime_api_trace && tool::get_config().benchmark_mode != tool::config::benchmark::sdk_callback_overhead) { + // Start each tool_init from a clean state: in attach/reattach + // scenarios the static-object maps persist across sessions and + // event pointers can be reused, so stale entries from a prior + // session would cause false-positive dependency resolutions. + tool::event_producer::clear_all(); + + ROCP_INFO << "Inter-stream HIP event-sync dependency tracking enabled"; + auto hip_event_dep_ctx = rocprofiler_context_id_t{0}; ROCPROFILER_CALL(rocprofiler_create_context(&hip_event_dep_ctx), diff --git a/projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp b/projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp index e1cf4564cec..a1a529f730c 100644 --- a/projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp +++ b/projects/rocprofiler-sdk/source/lib/tests/tool/event_producer_map.cpp @@ -54,8 +54,8 @@ sid(uint64_t v) class EventProducerMap : public ::testing::Test { protected: - void SetUp() override { ep::clear_for_testing(); } - void TearDown() override { ep::clear_for_testing(); } + void SetUp() override { ep::clear_all(); } + void TearDown() override { ep::clear_all(); } }; } // namespace @@ -163,7 +163,7 @@ TEST_F(EventProducerMap, stash_zero_corr_id_is_silent_noop) EXPECT_FALSE(ep::consume_resolved_wait(0).has_value()); } -TEST_F(EventProducerMap, clear_for_testing_resets_both_maps) +TEST_F(EventProducerMap, clear_all_resets_both_maps) { ep::record_event_producer(fake_event(0x600), sid(1), 1); ep::stash_resolved_wait(7, sid(2), 2); @@ -171,7 +171,7 @@ TEST_F(EventProducerMap, clear_for_testing_resets_both_maps) EXPECT_EQ(pre.producer_map_size, 1u); EXPECT_EQ(pre.resolved_wait_size, 1u); - ep::clear_for_testing(); + ep::clear_all(); auto post = ep::get_counters(); EXPECT_EQ(post.producer_map_size, 0u);