Repository navigation
Conversation
…ue callback registry Isolates the counter-collection slice of the callback-removal effort (reference PRs #8730 / #8586) into a small, single-service change, per review guidance to land callback removal one service at a time. Independent of kernel replay. Counter collection no longer registers a queue_callbacks_t with the HSA queue controller. Instead the HSA write interceptor and async signal handler call explicit free functions: - counters::write_hook / signal_completion_hook / is_any_active (new counters/queue_hooks.{hpp,cpp}) iterating active dispatch_counter_collection contexts and calling the existing queue_cb / completed_cb - hsa/queue_hooks/client_ids.hpp: stable inst_pkt producer tags The per-queue callback registry is intentionally retained for thread-trace, pc-sampling and spm; those migrate in follow-up PRs. The callback thread (callback_thread_start/stop) is preserved. The WriteInterceptor gate and should_batch predicate detect counters via is_any_active() (it no longer contributes to queue.get_notifiers()). Removes the now-unused queue_id from counter_callback_info and updates counter tests off iterate_callbacks. Co-authored-by: Cursor <cursoragent@cursor.com>
Address review feedback on #8891: completion routing must use registered contexts (packet_return_map provenance), stop_context must drain in-flight GPU work before deactivation, and queue-interposed services must stop while the context remains active for serialization coordination. Adds queue_hooks regression test for stop-then-complete ordering required by kernel replay (#7960/#8622). Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Reverts the queue_controller_sync() added earlier on this branch. The justification given for it was wrong: it claimed callback_thread_stop() could strand records belonging to dispatches still in flight. It cannot. - consumer_thread_t::exit() sets valid=false, notifies, then waits on the exited flag, and consumer_loop only sets exited once read_ptr == write_ptr, so everything already queued is consumed before the join. counters/tests/ consumer_test.cpp restart already covers this drain-on-exit path. - consumer_thread_t::add() falls back to calling consume_fn inline when valid is false, so a completion arriving after the thread is gone is still processed on the async signal handler thread. Nothing else in counters::stop_context needs an idle GPU. In-flight completions are handled by signal_completion_hook routing over registered rather than active contexts, and the serializer transition is reconciled GPU-side by the hsa_barrier that profiler_serializer::disable() pushes. This aligns the counters path with the same change on the SPM branch, prompted by review feedback from @SrirackshaNag noting that SPM was the only service with a sync in its stop path and that the hook migration removes its original reasons. The context::stop_context ordering is kept but re-justified: with no drain, the point is that disable_serialization() runs before dispatches stop being instrumented, so no dispatch is submitted without serializer packets while the serializer is still enabled. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Applies the changes requested in review of the base PR, so that merging this stacked PR clears them on users/vkale/remove-callbacks-counters. Renames, per @vlaindic's comment that 'write' does not convey that this is a write interceptor hook, and per the name agreed in the reply on that thread: counters::write_hook -> counters::kernel_dispatch_phase_enter_hook counters::signal_completion_hook -> counters::kernel_dispatch_phase_exit_hook The exit hook is renamed for symmetry even though only the enter hook was flagged; leaving one half renamed reads worse than either alternative, and the paired names make the enqueue-versus-completion asymmetry this branch introduces easier to follow. is_any_active keeps its name, being neither phase. SPM mirrors these names on its own branch and is deliberately left alone until the review of #8887 settles, so the two branches are briefly inconsistent by choice rather than by oversight. client_ids.hpp: the four constants become an unscoped enum with int64_t as its fixed underlying type. That answers the question asked in review -- the int64_t was not arbitrary, it matches hsa::ClientID, which is how the tag is stored in inst_pkt_t. Unscoped rather than enum class so the implicit conversion to ClientID survives and no emplace site needs a cast. The regression test now constructs the pair piecewise instead of through std::make_pair, which would deduce the tag's enum type and lean on a pair-to-pair conversion; piecewise converts the tag directly and matches the production hook. Copyright headers on the four files the base PR adds move from 2025 to 2026. Verified without a build, which is not available here: clang-format 11.1.0 (the version the repo's .clang-format pins) leaves no unrelated churn, no references to the old names remain, and a standalone translation unit confirms the enum converts to ClientID in both the piecewise and make_pair forms. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
… just unmapped Strengthens the stop-while-in-flight regression test to what the review on #8891 asked for: enqueue kernels, stop the context mid-flight, and assert every kernel produces a record. The previous version had two weaknesses. It passed nullptr for the record callback, so it could not observe delivery at all -- and delivery is the failure the review described. And it drove a single dispatch, which cannot distinguish 'routed the one entry present' from 'routed by provenance'. Now three dispatches are registered while the context is active, the context is stopped while all three are in flight, and completions are fired one at a time. Two independent invariants are checked, because they fail separately: - routing: the map is asserted to drop by exactly one per completion, not merely to reach zero at the end. An aggregate check would pass even if one completion erased another dispatch's entry. - delivery: every dispatch reaches the tool's record callback. Erasure happens in completed_cb before the hand-off to the consumer, so a map-only test would pass even if nothing were ever delivered. The delivery assertion also covers the lifetime condition raised in the review: callback_thread_stop() has already run by then, so these records arrive through the inline path in consumer_thread_t::add() rather than the consumer thread. No kernel executes, so counter values are whatever the profile's output buffer holds; the assertion is about delivery, not values. Not compiled here -- no ROCm toolchain in this environment. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…'s ClientID Captures the non-blocking observation from the review of #8790 at the declaration rather than only in a PR discussion, since that is where someone about to route on these values will look. The tags overlap the per-queue registry's auto-incrementing ClientID, which also starts at 1 and is still in use by the services not yet migrated. This is inert only because no consumer dispatches on the tag: counters' completed_cb and thread trace's post_kernel_call both identify their own packets by pointer lookup or dynamic_cast. The comment states what has to happen before that changes. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…ntext Restores the queue_controller_sync() this branch had earlier removed, per review priority: the review of #8891 accepted provenance-based completion routing on the condition that the callback thread and the counter_callback_info objects stay alive until in-flight dispatches drain, and the drain is what makes that hold literally rather than by argument about consumer_thread_t's fallback behaviour. An intermediate revision removed it on the reasoning that consumer_thread_t::exit() drains its queue and add() self-consumes inline afterwards. Both facts are true and they remain a useful backstop, but they are not the guarantee that was asked for -- they argue a late completion is still processed, rather than ensuring there is no late completion to process. The comment now says which is which. Ordering matters and is now stated at the call site: context::stop_context runs this function while the context is still in the active list, so for the duration of the drain the enter hook still reaches queue_cb, whose disabled path returns serialize=true and keeps the serialized->unserialized transition coordinated until disable_serialization() runs. That is what removes the need for a separate draining flag -- and it only works because the drain happens before the active slot is cleared. callback_thread_stop() now runs after the drain rather than alongside it. Not compiled here -- no ROCm toolchain in this environment. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…th drain The comment claimed provenance routing makes stop_context safe without a GPU drain, which contradicts counters::stop_context, where the drain was kept. State what each mechanism actually guarantees: provenance routing delivers a completion that arrives, the drain bounds when completions arrive. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Records how queue-interposed services are wired into dispatch interception after the per-queue callback registry is removed: why the enter hook routes over active contexts while the exit hook routes over registered ones, why that asymmetry is the difference between delivering and dropping the completion of an in-flight dispatch, and why the stop path drains the GPU. Folded in from a separate docs branch so it is reviewed alongside the code it describes. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
… its constructor
context::correlation_id declares constructors, deletes its copy and move, and
holds private ref counters, so it is not an aggregate and a designated
initializer does not compile:
error: initialization of non-aggregate type 'context::correlation_id'
with a designated initializer list
Default-construct and assign the one member the test needs, which leaves every
other member at the same default the designated initializer would have.
This broke the rocprofiler-sdk build, which is why all four Core test jobs and
Code Coverage failed alongside the TheRock package builds.
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…eue interception counters::stop_context() drains in-flight dispatches via queue_controller_sync() and coordinates serializer teardown, both of which require a live queue controller. finalize() ran queue_controller_fini() first and invoked the client finalizers afterwards, so stopping a counter-collection context during teardown operated on destroyed queue state and segfaulted. Also skip empty optional slots when walking the registered-context array; the array is sparse and dereferencing a disengaged optional is undefined. Fixes the hip-streams-per-thread SEGFAULT seen in CI. Co-authored-by: Cursor <cursoragent@cursor.com>
Moving invoke_client_finalizers() ahead of the subsystem teardowns deadlocked OMPT finalization: rocprofv3-test-ompt-tracing and rocprofv3-test-ompt-all timed out on every CI OS, and openmp-tools.test_retired_correlation_ids failed with them. The reorder also did not fix the hip-streams-per-thread SEGFAULT it was meant to address, so revert it and keep the original ordering. The sparse-optional guards in context.cpp are unaffected and stay in place. Co-authored-by: Cursor <cursoragent@cursor.com>
buffered-api-tracing and the two thread-trace-api-triple-buffer tests failed on ubuntu-22.04 only, while rhel-8.8, rhel-9.5 and sles-15.6 passed and TheRock's own rocprofiler-sdk shard passed. No content change; this re-runs CI on an identical tree to confirm. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Dispatch counter collection applied to the whole machine: a context had no notion of which GPUs it collected on, so serialization was enabled on every agent for as long as any counter-collection context was active, and only one such context could be active at all. A tool could decline to profile a dispatch by returning a null config, but that only suppressed instrumentation -- the dispatch was still serialized, which is the part that costs performance and the part the tool had no lever over. The public header has carried a note since the service was introduced saying the agent restriction was unresolved and left to the tool. Store the agent set on the context's counter-collection service and factor it into the three queries that need it: serialization scope at start/stop, the dispatch enter hook, and counter-collection context conflict. An empty set still means every agent, so existing tools are unaffected. Contexts with disjoint agent sets may now be active simultaneously, which is what makes "context A on GPU-0, context B on GPU-1" expressible. Serialization is now reference counted per agent. Counter collection, thread trace and SPM each enable it independently, and without a refcount whichever service stopped first silently unserialized the ones still running. That was a pre-existing bug; scoping serialization per agent makes it reachable more easily, so it is fixed here rather than left behind. Adds rocprofiler_dispatch_counting_service_set_agents() to declare the set. It is rejected while the context is active, since the dispatch path reads the set without a lock and start_context has already used it to scope serialization. Co-authored-by: Cursor <cursoragent@cursor.com>
The previous commit was formatted with clang-format 18 and left the source and cmake format checks failing. CI pins clang-format 11.1.0 and cmake-format 0.6.13, and v18 disagreed with v11 on two spots -- it joined a pre-existing ROCP_ERROR continuation in counters/core.cpp onto one line and moved the extern "C" brace in dispatch_counting_service.cpp -- neither of which this change had any reason to touch. Co-authored-by: Cursor <cursoragent@cursor.com>
…ware The unit tests assert the internal state, which cannot show the property a user actually cares about: that a GPU outside the context's agent set keeps running its streams concurrently. This adds an integration test that measures a multi-stream workload on GPU-0 while a counter collection context is scoped to GPU-1, and fails if GPU-0 gets serialized or instrumented. Machines differ too much for an absolute time to mean anything, so every timing assertion is a ratio against a baseline measured in the same process, and the unrestricted context serves as a positive control -- when it does not visibly serialize the same workload the machine cannot demonstrate stream concurrency and the timing assertions are skipped instead of guessed at. The record and callback assertions are exact and always run. The test reports SKIP and exits 0 below two GPUs. Verified against a deliberately reverted build: without the agent filtering the scoped GPU-0 workload goes from 7.7ms to 32.0ms and three checks fail. Co-authored-by: Cursor <cursoragent@cursor.com>
Suggested by @jrmadsen while reviewing #9868, applied as written: hsa/queue_controller.hpp using agent_handle_set_t = std::unordered_set<rocprofiler_agent_id_t> context/context.hpp std::unordered_set<rocprofiler_agent_id_t> agents{} context/context.hpp return agents.empty() || agents.count(agent_id) > 0 Also pulls in <rocprofiler-sdk/cxx/hash.hpp> and <rocprofiler-sdk/cxx/operators.hpp>, which he noted give full support for using the opaque handles in map/set/uomap/etc. Those includes are what let the sets key on rocprofiler_agent_id_t directly instead of unwrapping to the uint64_t handle. The cascading retyping this implies for the serialization refcounts and their call sites follows in the next commit, so that this one holds only the suggested text. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…ects out of line Follows the suggestions applied in the previous commit. Retyping the two agent sets to rocprofiler_agent_id_t leaves the surrounding code still speaking in raw uint64_t handles, so this carries the change through: - serialization_refcount::per_agent keys on rocprofiler_agent_id_t, and count(), enabled() and any() take/iterate the typed id - queue_controller.cpp and counters/core.cpp stop unwrapping to ::handle at the call sites - comments that described the sets as holding rocprofiler_agent_id_t::handle are reworded, since they no longer do Also addresses jrmadsen's fourth review thread, which asked for the body to live in the .cpp: dispatch_counter_collection_service::intersects is now declared in context.hpp and defined in context.cpp. He described that one rather than writing it, so it sits here rather than in his commit. Co-authored-by: Cursor <cursoragent@cursor.com>
Use rocprofiler_agent_id_t values in per_agent_scoping_test serialization calls instead of raw handles, and apply clang-format 11 to queue_controller. Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
documentation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
There was a problem hiding this comment.
Pull request overview
Integration preview branch composing multiple in-flight PRs to (1) remove per-queue callback-registry wiring in favor of explicit per-service queue hooks, and (2) introduce experimental kernel replay (multi-pass per-dispatch execution with snapshot/restore) for counter collection, including rocprofv3 CLI/docs and regression tests validating replay behavior and concurrency fixes.
Changes:
- Adds experimental kernel replay callback-tracing domain, memory tracker/snapshot/restore, and localized per-pass context override mechanism.
- Migrates counters, thread trace, PC sampling, and SPM off the per-queue callback registry via explicit
*/queue_hooks.{hpp,cpp}called fromhsa/queue.cpp. - Extends rocprofv3 and output (CSV/JSON) to support kernel replay, and adds new integration/unit tests (including deterministic concurrency regression coverage).
Reviewed changes
Copilot reviewed 84 out of 88 changed files in this pull request and generated 2 comments.
Show a summary per file
| File | Description |
|---|---|
| projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/pytest.ini | Adds pytest configuration for kernel replay validation. |
| projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/conftest.py | Adds pytest CLI options/fixtures for validating replay JSON output structure. |
| projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/CMakeLists.txt | Adds ctest generate+validate coverage for rocprofv3 kernel replay output and interval regression. |
| projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt | Registers the rocprofv3 kernel replay test directory. |
| projects/rocprofiler-sdk/tests/kernel-replay-concurrency/main.cpp | Adds deterministic workload for kernel replay concurrency regression testing. |
| projects/rocprofiler-sdk/tests/kernel-replay-concurrency/CMakeLists.txt | Builds and registers deterministic kernel replay concurrency integration test. |
| projects/rocprofiler-sdk/tests/kernel-replay-concurrency/client.cpp | Adds LD_PRELOAD tool driving replay/victim opt-out and handshake for deterministic race reproduction. |
| projects/rocprofiler-sdk/tests/CMakeLists.txt | Registers kernel-replay-concurrency tests in the overall test suite. |
| projects/rocprofiler-sdk/tests/bin/kernel-replay/CMakeLists.txt | Adds HIP test binary used by rocprofv3 kernel replay integration tests. |
| projects/rocprofiler-sdk/tests/bin/CMakeLists.txt | Registers the kernel-replay test binary. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/queue_hooks_test.cpp | Adds unit test for thread-trace queue hook activeness predicate. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/CMakeLists.txt | Builds/registers thread-trace queue-hooks unit test. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/queue_hooks.hpp | Introduces explicit thread-trace write/completion hooks and activeness check. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/queue_hooks.cpp | Implements thread-trace explicit queue hooks. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.cpp | Removes queue-controller registration and adds replay-local override gating in pre-kernel hook. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt | Adds thread-trace queue hook sources/headers to build. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/tests/queue_hooks_test.cpp | Adds unit test for SPM queue hook activeness predicate. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/tests/core.cpp | Updates SPM unit tests to stop asserting callback-registry state. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/tests/CMakeLists.txt | Adds SPM queue hook test source. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/queue_hooks.hpp | Introduces explicit SPM write/completion hooks and activeness check. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/queue_hooks.cpp | Implements SPM explicit queue hooks. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/core.hpp | Removes stored per-queue callback ClientID from SPM callback info. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/core.cpp | Removes queue-controller callback registration/removal; relies on explicit hooks. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/spm/CMakeLists.txt | Adds SPM queue hook sources/headers to build. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp | Installs kernel replay memory tracker wrappers and includes replay in queue-interposition gating. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/types.hpp | Removes unused queue interceptor ClientID field. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/tests/queue_hooks_test.cpp | Adds unit test for PC sampling agent-configuration predicate. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/tests/CMakeLists.txt | Adds PC sampling queue hook test source. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/queue_hooks.hpp | Introduces always-linkable PC sampling hooks/predicates. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/queue_hooks.cpp | Implements always-linkable PC sampling predicates and explicit completion hook. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.hpp | Exposes PC sampling kernel completion entry point for explicit hook call. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp | Removes callback registration; routes completion via explicit hook. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/CMakeLists.txt | Ensures queue_hooks builds even when PC sampling HSA support is unavailable. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/utils.hpp | Adds shared allocation-classification helper API for tracker/tests. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/utils.cpp | Implements allocation classification via pointer-info query. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/snap_kernels.hpp | Declares host-callable wrappers around HIP test kernels for snapshot/restore tests. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/snap_kernels.cpp | Implements HIP kernels used by snapshot/restore tests. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/memory_tracker_test.cpp | Adds host-only tests for tracker install gating and inventory behavior. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/memory_snapshot_test.cpp | Adds host-only tests for snapshot ok/empty/restore gating contract. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/local_context_test.cpp | Adds unit tests for replay-local per-pass context override semantics. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/CMakeLists.txt | Builds/registers kernel replay unit tests and manages GPU resource locking. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/replay_callbacks.hpp | Declares kernel replay callback execution and replay-plan structures. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_tracker.hpp | Declares replay memory allocation inventory and HSA table wrapper installers. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_tracker.cpp | Implements allocation tracking wrappers, inventory, and executable/kernarg exclusions. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.hpp | Declares minimal device memory snapshot/restore structures and contract. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/local_context.hpp | Declares replay-local per-pass context override mechanism. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/local_context.cpp | Implements thread-local loop-scoped override map and arm-window enforcement. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/kernel_replay.hpp | Declares enum-to-string utilities for kernel replay operations. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/kernel_replay.cpp | Implements enum-to-string utilities for kernel replay operations. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/CMakeLists.txt | Adds kernel replay sources/headers and tests to build. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp | Changes Queue::sync() API to return a drain-success boolean for replay/teardown behavior. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/CMakeLists.txt | Adds header-only client-id tags build target wiring. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp | Introduces stable inst-packet producer tags for explicit hooks. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp | Ensures queue interposition stays enabled when kernel replay tracing is configured. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt | Registers the new hsa/queue_hooks subdirectory. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp | Updates counter tests to stop relying on callback registry state. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt | Adds counters queue hook test source. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.hpp | Introduces explicit counters enter/exit hooks and activeness check. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.cpp | Implements explicit counters queue hooks, including registered-context completion routing. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp | Adds replay-local context override gating into counters dispatch path. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp | Removes stored per-queue callback ClientID from counters callback info. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp | Removes per-queue callback registration/removal and tightens drain/teardown ordering. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/CMakeLists.txt | Adds counters queue hook sources/headers to build. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp | Stops queue-interposed services before clearing active slot (serialization ordering fix). |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/CMakeLists.txt | Registers new kernel_replay subdirectory in SDK library build. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp | Adds KERNEL_REPLAY callback-tracing kind integration and enables tracker/replay service gate. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp | Adds kernel replay config flag (currently conflicting with base output_config member). |
| projects/rocprofiler-sdk/source/lib/output/output_config.hpp | Adds kernel replay field to output config serialization. |
| projects/rocprofiler-sdk/source/lib/output/output_config.cpp | Parses ROCPROF_KERNEL_REPLAY into output config. |
| projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp | Adds optional Replay_Pass column emission and replay-specific CSV encoder. |
| projects/rocprofiler-sdk/source/lib/output/csv.hpp | Adds counter_collection_replay_csv_encoder type alias. |
| projects/rocprofiler-sdk/source/lib/output/counter_info.hpp | Adds per-record replay pass index to counter records and serialization. |
| projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h | Adds KERNEL_REPLAY callback-tracing kind and kernel replay operation enum. |
| projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/kernel_replay.h | Adds experimental public API payload struct for kernel replay callbacks and localized context control. |
| projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/CMakeLists.txt | Installs the new experimental kernel replay header. |
| projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp | Adds enum-string metadata and static asserts for new kernel replay kind/operations. |
| projects/rocprofiler-sdk/source/docs/how-to/using-rocprofv3.rst | Documents rocprofv3 kernel replay beta flag behavior and constraints. |
| projects/rocprofiler-sdk/source/docs/conceptual/queue_hooks/queue_callback_removal_architecture.md | Adds conceptual architecture doc for callback registry removal and routing/teardown invariants. |
| projects/rocprofiler-sdk/source/bin/rocprofv3.py | Adds --kernel-replay-beta-enabled CLI flag, validation, and env/config merging logic. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+158
to
+162
| // Route counter collection through in-process kernel replay: collect every --pmc counter group | ||
| // in a single application run by replaying each dispatch once per group (device-memory | ||
| // snapshot/restore between passes). The pass count is the number of counter groups. | ||
| bool kernel_replay = get_env("ROCPROF_KERNEL_REPLAY", false); | ||
|
|
Contributor
Author
There was a problem hiding this comment.
@copilot Fix the code for all comments in this review comment.
When a review comment includes a suggested change, apply the suggestion exactly.
Do not make changes beyond what is described in the linked review comment.
Contributor
There was a problem hiding this comment.
Fixed in the latest commit. Removed kernel_replay from the derived config struct (lines 158–162) and its CFG_SERIALIZE_MEMBER entry from config::save(). The base output_config already declares the field (initialized to false), already reads it from ROCPROF_KERNEL_REPLAY in parse_env() (called by load_from_env() which the config constructor invokes), and already serializes it — so there is now a single source of truth.
Comment on lines
+182
to
+183
| plan.user_data = plan.config_contexts.empty() ? tracing::empty_user_data | ||
| : plan.config_contexts.front().user_data; |
Contributor
Author
There was a problem hiding this comment.
@copilot Fix the code for all comments in this review comment.
When a review comment includes a suggested change, apply the suggestion exactly.
Do not make changes beyond what is described in the linked review comment.
Please add pertinent tests and comment changes.
Contributor
There was a problem hiding this comment.
Fixed in the latest commit. Changed config_contexts.front() to config_contexts.back() so that user_data is taken from the same last-registered context as pass_count_cb (which comes from the shared config_data payload under last-writer-wins semantics). Added an explanatory comment at the fix site in replay_callbacks.cpp. Added replay_plan_test.cpp with three host-only unit tests (single_context_uses_that_context, multiple_contexts_uses_last_not_first, empty_contexts_gives_empty_user_data) that validate the selection invariant directly on callback_context_data_vec_t without requiring the HSA or rocprofiler runtime; registered the new target kernel-replay-plan-test in CMakeLists.txt.
Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
…r matches 5072c17 skipped module variables whose HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SEGMENT is READONLY. ROCr derives that value from SHF_AMDGPU_HSA_READONLY (Symbol::Segment() in libamdhsacode/amd_hsa_code.cpp), a legacy HSA code object section flag that LLVM does not emit. Compiling a HIP file for gfx942 with clang 18 puts a __constant__ int in .data (WRITE, ALLOC) next to the __device__ globals, and no section carries 0x00200000, so every variable reports GLOBAL and the filter skipped nothing. Keep capturing __constant__ variables, and record why neither IS_CONST nor VARIABLE_SEGMENT can single them out. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…constant comment Where __constant__ variables land depends on the compiler. Ubuntu clang 18 puts a __constant__ int in .data, but the amdclang++ 23 that CI builds with puts it in a non-writable section: on #12605 the IS_CONST filter left the excludes_read_only_module_constant fixture in the snapshot, which only happens when its section lacks SHF_WRITE. Neither placement sets SHF_AMDGPU_HSA_READONLY, so the comment's reasons for not filtering on IS_CONST or VARIABLE_SEGMENT hold either way. State those reasons without the placement claim. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…l handler thread Completed thread trace packets no longer hold a reference to their ThreadTracerAgent. If a traced dispatch outlived finalize's queue sync, the packet held the last reference, so the agent was destroyed on ROCr's async signal handler thread. On the HSA backend its destructor calls hsa_queue_destroy(), which waits for a handler that only that thread runs, so the process hung at exit. post_kernel_call goes back to develop's order: decrement the pending count, then look the agent up under agents_map_mut. A completion that arrives after resource_deinit() drops its trace, as on develop. resource_deinit() now destroys the agents after releasing the lock, because the exit hook can be waiting for it on the async signal handler thread. Also: - Make start_context()/stop_context() non-const again and drop the mutable on enabled; neither was needed. - Return early from the exit hook when a completion carries no thread trace packets, as SPM's hook does. - Delete enter_hook_ignores_dispatches_on_other_agents, which could not fail because the agent map was never built. - Correct test and code comments that overstated what they check or said other services still use the registry. Co-authored-by: Cursor <cursoragent@cursor.com>
…o the kernel replay integration Brings in #11967's review fix f5eb45a: completed thread trace packets no longer keep their ThreadTracerAgent alive, so an agent is never destroyed on ROCr's async signal handler thread, where on the HSA backend hsa_queue_destroy() would wait on itself and hang the process at exit. With all four services migrated, nothing calls QueueController::add_callback here, so the merge also drops the parentheticals at the counter collection, thread trace and SPM enter hooks in hsa/queue.cpp that said other services still flow through signal_callback. Co-authored-by: Cursor <cursoragent@cursor.com>
…PM start The start-transition fix that came with #11968 (cb5bf4e) closes three windows in start_context(), and its two tests cover two of them: a stop landing before the SPM service starts, and two overlapping starts admitted onto one agent. The third is rocprofiler_spm_dispatch_counting_service_ set_agents() replacing the agent set of a context whose start has passed its conflict scan but not yet published its slot, which starts the context on an agent it was never checked against, or frees the set while enable_serialization() iterates it. set_agents_does_not_race_context_start holds agent 0 with one context and starts a second, restricted to agent 1, 1000 times while another thread flips that context between the two agents. Every start must either conflict or come up on agent 1 only. On develop it passes 20 of 20 runs. With cb5bf4e's context.cpp hunks reverted it fails 11 of 20 runs, all with a segmentation fault (the two existing tests fail 20 of 20). Co-authored-by: Cursor <cursoragent@cursor.com>
start_context() marks a context as start-in-progress until its services are running, so a concurrent stop or start waits rather than acting on a half-started context. It releases that marker before start_agent_ctx(), which calls the tool's profile callback synchronously, and PC sampling was started after that release too. A stop that claimed the context in the gap found the service not yet enabled and left it alone, cleared the active slot, and the start then enabled PC sampling on a context that was no longer active: nothing could stop it, and the next start of that context failed with ROCPROFILER_STATUS_ERROR because the service was already enabled. Start PC sampling before the marker is released. It calls no tool callback, so it cannot re-enter the lifecycle API while holding the marker, and it never shares a context with counter collection, so its status cannot race start_agent_ctx()'s. pc_sampling.concurrent_start_stop_does_not_leave_sampling_enabled starts and stops one PC sampling context from two threads, 20000 times each. On MI325X it fails 9 of 40 runs without this change (each failure a start returning ROCPROFILER_STATUS_ERROR) and passes 20 of 20 with it; counting earlier versions of the test with fewer iterations, 41 of 220 runs failed without the change and 0 of 110 with it. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
Same change as ac9f915 on the #11970 docs branch: the phase-two comment in context::stop_context now names the teardowns that wait on the GPU (counter collection's and SPM's call hsa::queue_controller_sync(), device thread trace waits for its stop packets, dispatch thread trace does not wait), and the counter collection comments and the interposition_sync() header say that only the per-queue part of queue_controller_sync() is bounded; the completion-monitor drain before it has no time limit. Comments only.
Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
develop now carries the final thread trace (#11967) and counters (#11970) callback removals. Conflicts in context/context.cpp, counters/core.cpp, counters/sample_processing.cpp and hsa/queue.cpp were resolved as unions: - Git applied both develop's and this branch's insertion of the same code at different positions, so one copy of each was dropped: the counters and thread trace enter and exit hook calls in hsa/queue.cpp (this branch's placement kept), and dispatch_counter_collection_service::intersects() in context/context.cpp. Every service hook is called once. - start_context() keeps this branch's PC sampling start under the context start marker (d4b7863) and takes develop's new device SPM start (#11879) after the counters agent start. - The interception and per-packet gates take develop's form, which ANDs or ORs the same per-agent predicates; comments keep this branch's wording from c590486. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
cursor
Bot
changed the title
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…eplay integration Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…y integration Co-authored-by: Vivek Kale <vivek.kale@amd.com>
… integration Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…play integration Co-authored-by: Vivek Kale <vivek.kale@amd.com>
cursor
Bot
changed the title
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
CI's black --required-version 26 collapses a few call-argument wraps that black 25 left alone. Matches the python (3.10) check on the integration preview. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
… the kernel replay integration Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Co-authored-by: Vivek Kale <vivek.kale@amd.com>
This branch has not been deployed
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.
What this is
An integration preview, not a merge candidate. It composes
developwith the current heads ofthe four callback-removal PRs so that the kernel replay service can be exercised against all four
migrated subsystems at once. Each dependency is reviewed and merged on its own PR; this branch
exists only to show that they compose and that replay still works on top of the composition.
This must not merge, and must not merge ahead of its dependencies. Apart from merge-conflict
resolutions it carries no content of its own.
Composition
Merged onto
develop, in this order:users/vkale/remove-callbacks-counters--with-cursor-fixesusers/vkale/remove-callbacks-thread-trace--with-cursor-fixesusers/vkale/remove-callbacks-pc-sampling--with-cursor-fixesusers/vkale/remove-callbacks-spm--with-cursor-fixesThe kernel replay service is no longer carried on this branch. #7960 and #10439 are already merged
into
develop, so replay arrives with the base and is now newer here than the version this branchused to carry.
Why it was recomposed rather than rebased
The previous revision merged the older fork-hosted PRs (#8891, #8790, #8895, #8887) against an
August
develop. Since then the thread trace and SPM work merged into their--with-cursor-fixesbranches, those branches were rebased onto a later
develop, and the kernel replay service landedin
develop. The merge bases this branch relied on no longer exist, so it was rebuilt from thecurrent heads instead of rebased.
#8891 and #8895 are deliberately left untouched as review-history records. Their live equivalents
are #11970 and #11969.
Conflict resolutions
Every conflict was the same shape: two subsystems independently migrating the same call site off
the per-queue callback registry. All were resolved as unions, never by taking one side.
hsa/queue.cpp— every subsystem's explicit enter and exit hook is invoked at the call sites theregistry used to drive. The interception gate and the per-packet-mode gate AND together every
subsystem's per-agent predicate, so a run using only one of them still intercepts, and queues on
agents no service is configured for stay on the fast path.
context/context.cpp— the context-registry snapshot rework appears exactly once; it is identicalon the three branches that carry it. The
start_contextconflict cascade keeps each service's ownagent-disjointness clause rather than the unconditional check it replaced.
thread_trace/core.cpp—resource_deinit()andpost_kernel_call()take the migrated threadtrace versions, so a completion that races teardown is delivered through the reference the packet
carries instead of a lookup in a map that teardown may already have cleared.
What the committed tests assert
ThreadTraceQueueHooks.CompletionAfterResourceDeinitStillDrainsasserts that a thread tracecompletion arriving after
resource_deinit()is still drained rather than dropped.unit.counters_per_agent.disjoint_agent_contexts_can_be_active_togetherandoverlapping_agent_contexts_still_conflict, with the equivalent thread trace and SPM cases,assert that contexts scoped to disjoint agent sets coexist while overlapping ones still conflict.
test-kernel-replay-local-context-*integration tests exercise replay against countercollection, thread trace, SPM and PC sampling, including stop-and-restart sequences.
rocprofv3-test-kernel-replay-validateasserts the replay pass column covers every pass exactlyonce per
dispatch_id.The SPM integration tests remain
set(IS_DISABLED True)per the SPM owner's decision; that file isuntouched on this branch.
What this is
An integration preview, not a merge candidate. It composes
developwith the open kernel-replay product PRs so they can be exercised together: range replay, the kernel-replay performance stack, the benchmark harness, and the beta-flag gating. Each dependency is reviewed and merged on its own PR; this branch exists only to show that they compose.This must not merge, and must not merge ahead of its dependencies. Apart from merge-conflict resolutions it carries no content of its own.
Benchmark harness (requested link)
users/vkale/kr-perf-0-replay-benchmark-harness--with-cursor-helpComposition (current head)
Merged onto
develop, in this order (all clean, or a trivial changelog/docs union):users/vkale/range-replay-v1__hip_cuid_*, retained staging)users/vkale/kr-perf-5-retained-staging--with-cursor-helpusers/vkale/kr-perf-0-replay-benchmark-harness--with-cursor-helpusers/vkale/kr-perf-1-skip-constant-symbols--with-cursor-help--kernel-replay-beta-enabledrequires--replay-mode kernelusers/vkale/kernel-replay-beta-flag-with-cursor-fixes-helpNot on this preview (conflicts; left out)
hsa/queue.cppandkernel_replay/memory_snapshot.{hpp,cpp}. Both rewrite the same snapshot/restore path: this preview carries retained host staging + sync host↔device restore + phase timings; #10916 replaces that with GPU-local/pinned arenas, a richersnap(capture_context_t), and async blit restore submitted ahead of the next pass. Not a clear union — picking one side drops the other stack's design. Needs a dedicated compose/rebase before it can land here.CHANGELOG.md, docs TOC/index,rocpd/features.py, andversions.yml. Needs a dedicated rebase onto current develop / this tip.Other open kernel-replay-adjacent PRs that are not part of this product composition (experiments, rocprofv3 feature work, or other stacks): #12606 (todos on top of this preview), #13083 / #13168 (pass-toggle experiments), #12776 / #12729 (ATT/json pass tooling), #12888 (serial test runner), #13164 (rocpd repeated PMC), #12468 (overnight stress), #12815 / #13272 (weekly range-replay test stack).
Conflict resolutions on this refresh
CHANGELOG.mdwhen merging perf(rocprofiler-sdk): keep kernel replay snapshot staging across dispatches #12610 — both sides added### Changedentries; kept both.docs/conceptual/kernel_replay/index.mdwhen merging test(rocprofiler-sdk): add a kernel replay benchmark harness and benchmarking guide #12885 — kept the range-replay row from this branch and the benchmarking paragraph / more precise counter-group-rotation wording from test(rocprofiler-sdk): add a kernel replay benchmark harness and benchmarking guide #12885.No other conflict resolutions.
#10916and#11422were aborted rather than force-fitted.Sync cadence
An every-other-day timer merges
developand the current heads of the carried PRs above into this branch (fast-forward push only). New conflicts abort the offending merge and get reported instead of guessed.What the committed tests assert
--replay-mode kernelgating from fix(rocprofiler-sdk): reject --kernel-replay-beta-enabled without --replay-mode kernel #11432.