rocprofv3: surface inter-stream HIP event-sync dependencies on wait/sync records - #1
Draft
ajassani wants to merge 3 commits into
Draft
rocprofv3: surface inter-stream HIP event-sync dependencies on wait/sync records#1ajassani wants to merge 3 commits into
ajassani wants to merge 3 commits into
Conversation
…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.
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.
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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
rocprofv3 already shows the per-record stream ID on
hipStreamWaitEventandhipEventSynchronizetool records, but it does not yet surface which stream and whichhipEventRecordcall the waiter is actually depending on. As a result, tools that want to reconstruct a cross-stream / critical-path dependency graph from a captured trace have to re-run the workload under a separate analysis pass or guess at the topology from timestamps.This PR closes that gap by adding two new fields on the HIP API tool record:
wait_on_stream— the stream that recorded the awaitedhipEvent_twait_on_correlation_id— the rocprofv3 correlation id of thathipEventRecordcallBoth default to
0(the "no resolved dependency" sentinel) when the awaited event was never observed being recorded earlier in the same trace, so the change is non-breaking from a record-shape standpoint.Technical Details
Event-lifecycle tracking utility
A new utility in
source/lib/rocprofiler-sdk-tool/event_producer_map.{hpp,cpp}maintains two process-global, shared/unique-locked maps using the existingcommon::Synchronized+common::static_objectbuilding blocks:hipEvent_t -> { stream_id, hipEventRecord corr_id }. Written onPHASE_EXITofhipEventRecord(and_spt/WithFlagsvariants); erased onhipEventDestroy.waiter corr_id -> { wait_on_stream, wait_on_corr_id }. Written onPHASE_EXITofhipStreamWaitEvent(and_spt) andhipEventSynchronizeafter a producer lookup; drained at buffer-flush time when the corresponding tool record is materialized.Callback wiring
A dedicated callback context is registered on
ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_APIfiltered to just the seven operation IDs involved in event lifecycle and synchronization. The context is gated onhip_runtime_api_traceso there is zero overhead when HIP tracing is off. The callback runs onPHASE_EXITonly (after the underlying HIP call has completed) so the producer state is consistent with what other threads will observe via lookup, and the body is wrapped in a top-leveltry/catch(the underlyingstd::unordered_mapops can throwstd::bad_allocunder memory pressure and this code runs synchronously inside the SDK callback dispatch path).Buffer-flush integration
At the existing
buffered_tracing_callbackflush site, the materialization oftool_buffer_tracing_hip_api_ext_record_tconsumes any stashed resolved-wait entry for the record's correlation id and routes through an extended constructor that fills the new fields.The lookup is gated on the record's operation kind — only
hipStreamWaitEvent,hipStreamWaitEvent_spt, andhipEventSynchronizecan ever carry a resolved dependency in the side-table, so the (locking) map call is skipped for the other ~99.95% of HIP API records on a representative DDP workload.Reattach hygiene
Both maps live in
common::static_objectand persist acrosstool_init→tool_detach→tool_initcycles.hipEvent_tpointers can be reused across sessions, so we clear both maps attool_initright before the dependency context is created. A singleROCP_INFOis also emitted at the same site so users see the service being enabled in the tool init log.Output coverage
out_hip_api_trace.csv) widens from 7 to 10 columns: the existing 7 +Stream_Id+Wait_On_Stream_Id+Wait_On_Correlation_Id. A dedicatedhip_api_csv_encoderis used only for HIP records so other CSV outputs are unchanged.out_results.json) picks upwait_on_stream(as{"handle": N}) andwait_on_correlation_idonhip_apibuffer records via cereal serialization. No schema breakage; the fields default to0.args(wait_on_stream_ID,wait_on_corr_id) on the relevant slices and, whenwait_on_correlation_id != 0, adds a secondFlow::ProcessScopedterminating at the resolved producer so the cross-stream dependency shows up as an arrow in the Perfetto UI.Deferred to follow-up PRs
Both formatters continue to emit valid records today; they just don't yet carry the new fields.
The HIP API CSV (
<out>_hip_api_trace.csv) now has 10 columns instead of 7. Consumers that parse this CSV positionally or hard-code the column count will need to be updated.The change is contained: only HIP records use the new
hip_api_csv_encoder<10>; other domains (kernels, markers, memory, etc.) keep their existing 7-column shape via the originalapi_csv_encoder<7>.The added columns are:
Stream_IdWait_On_Stream_Id0if no resolved dependencyWait_On_Correlation_IdhipEventRecordcorrelation id —0if no resolved dependencyIf a downstream gating flag for the new columns is preferred, happy to add an opt-in option (
--hip-trace-depsor similar) so the default shape stays at 7 columns until consumers have migrated.The JSON and Perfetto outputs are additive: new fields appear on existing records but no existing fields are renamed, reshaped, or removed.
JIRA ID
N/A — external contributor (no SWDEV access). Happy to have an internal champion attach a ticket if that helps with the merge gate.
Test Plan
Unit tests
source/lib/tests/tool/event_producer_map.cpp— 13 cases against the newtool-testsbinary (a focused gtest unit-test executable for the tool-side state machine):lookup_unknown_event_returns_nulloptlookup_null_event_returns_nulloptrecord_then_lookup_round_tripsrecord_overwrites_prior_entryrecord_null_event_is_silent_noopforget_event_removes_entryhipEventDestroycleanupforget_unknown_event_is_silent_noopdistinct_events_are_independentstash_then_consume_resolved_wait_drainsconsume_unknown_corr_id_returns_nulloptstash_zero_corr_id_is_silent_noop0is reserved as the no-dep sentinelclear_all_resets_both_mapsconcurrent_record_and_forget_does_not_corruptIntegration tests
tests/rocprofv3/hip-interstream-sync/— 1 execute + 5 validate cases. A small HIP workload exercises a cross-streamhipStreamWaitEventplus a hosthipEventSynchronize, then the validator asserts:hipStreamWaitEventandhipEventSynchronizerow carries a non-zeroWait_On_Correlation_Idwhose value matches some priorhipEventRecordrow'sCorrelation_Id(no dangling references);Wait_On_Stream_Iddistinct from their ownStream_Id(cross-stream is the entire point);0sentinel (no leakage onto records that shouldn't carry a dependency);Local validation
tool-testsgtest binary: 13/13 pass on an MI300X / MI250 debug build.ctest -R hip-interstream-sync: 6/6 pass (1 execute + 5 validate).rocprofv3 --hip-trace --kernel-trace --output-format json csv pftraceon a 2-rank PyTorch DDP workload: JSON, CSV, and pftrace all show non-zerowait_on_correlation_idfor the cross-stream waits; reconstructed dependency DAG matches the expectedbackward → reducetopology (78 of 149,872 records carry resolved metadata, ≈0.05%).Test Result
Submission Checklist
integration-testslabel, runs under existing GPU-required test infrastructure).Draft status
This PR is open on the fork (
ajassani/rocm-systems) for review and CI iteration before being promoted to an upstream PR againstROCm/rocm-systems. The OTF2 and Rocpd follow-ups will gate the upstream submission.