refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry by amd-vkale · Pull Request #11967 · ROCm/rocm-systems

amd-vkale · 2026-09-21T20:03:16Z

Base PR of the #8790 stack. Same content as #8790; the head branch users/vkale/remove-callbacks-thread-trace--with-cursor-fixes
lives in ROCm/rocm-systems rather than a fork, so CodeQL can read the container
registry secrets and the full CI matrix actually runs. #8790 keeps the review history
and is now stacked on this branch.

Motivation

This PR isolates the thread-trace 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.

Underlying Problem

Dispatch thread trace (ATT) currently registers a queue_callbacks_t on every HSA queue. That hides “is this service active?” inside a per-queue map, makes enqueue and completion the same walk, and serializes every GPU while any dispatch ATT context is active.

This PR:

Stops registering dispatch ATT with that map. The write interceptor and async signal handler call explicit free functions instead (thread_trace::write_hook, signal_completion_hook, is_any_active).
Scopes dispatch ATT to configured GPU agents so disjoint agent sets can be active together; serialization is refcounted per agent and overlapping contexts conflict via intersects().

Device thread trace is unchanged in its hook wiring; only dispatch ATT migrates here. The per-queue callback registry remains for counters, PC sampling, and SPM until those services land in their own PRs.

Notes on Bigger Picture

Sibling migrations: counter collection (#8891), SPM (#8887), PC sampling (#8895).

Kernel replay (#7960, #8622) benefits from callback removal but this PR is independent and behavior-preserving on its own.

Technical Details

The two hooks route differently, and that asymmetry is deliberate. write_hook iterates
active contexts, because injecting an ATT packet is only correct while the context is running.
signal_completion_hook iterates registered contexts, because a dispatch that was tagged
before stop_context still has to be collected afterwards. post_kernel_call self-filters by
THREAD_TRACE_CLIENT_ID and then dynamic_casts to TraceControlAQLPacket, so routing by
provenance is safe and differs from the counter/SPM packet_return_map mechanism.

flowchart LR
  WriteInterceptor --> WriteHook["thread_trace::write_hook: ACTIVE contexts + collects_on"]
  AsyncSignalHandler --> CompletionHook["signal_completion_hook: REGISTERED contexts, filtered by packet tag"]
  StartContext --> Serialize["enable_serialization(configured agents)"]
  StopContext --> Unserialize["disable_serialization(configured agents)"]

Permalinks pin to fe5f483.

Hooks and call sites

Declarations: queue_hooks.hpp#L42-L65
Enter — active contexts, collects_on, THREAD_TRACE_CLIENT_ID: queue_hooks.cpp#L45-L70
Exit — registered contexts, filtered by packet tag: queue_hooks.cpp#L72-L89
is_any_active: queue_hooks.cpp#L91-L95
no_real_consumers gate: queue.cpp#L441-L446
Enter call site: queue.cpp#L842-L852
Exit call site: queue.cpp#L298-L305
Batching disabled while active: queue.cpp#L1222-L1223
start_context / stop_context (no add_callback; serialization only): core.cpp#L597-L620
collects_on / intersects: core.cpp#L571-L585
post_kernel_call provenance filter: core.cpp#L531-L545
Context conflict uses intersects: context.cpp#L311-L317
Producer tags: client_ids.hpp#L50-L60

Per-agent scoping is configured via rocprofiler_configure_dispatch_thread_trace_service agent params (not a separate set_agents API). Agents are fixed at configure time.

Per-agent serialization is reference counted

Counters, SPM and ATT each ask for serialization independently. Once ATT can scope itself to a
subset of agents, a plain enable/disable flag is no longer sufficient — whichever service stopped
first would silently unserialize the ones still running. QueueController therefore keeps a
per-agent refcount, with a separate all counter for unscoped callers so that an unscoped request
still covers agents whose serializer has not been created yet.

Refcount state and enabled() / any(): queue_controller.hpp#L155-L183
update_serialization — only serializers whose effective count crossed zero are toggled, and the applied state is derived from the post-update count rather than the caller's flag: queue_controller.cpp#L727-L795
serializer() holds the refcount read lock across lookup/insert so a lazily created serializer cannot be born in the wrong state: queue_controller.cpp#L676-L712

Lock order is refcount-write → serializer-write in update_serialization and refcount-read →
serializer-upgrade in serializer(), i.e. the same order in both, so the pair is deadlock-free.

Reviewer Guide

Confirm the enter/exit asymmetry is what you expect: write_hook uses get_active_contexts (L58) while signal_completion_hook uses get_registered_contexts (L83). This is what makes in-flight completion safe after stop.
ATT-only runs still enter WriteInterceptor via thread_trace::is_any_active() (queue.cpp#L441-L446).
DispatchThreadTracer::start_context does not call add_callback (core.cpp#L597-L607).
context::stop_context now stops queue-interposed services before clearing the active slot (context.cpp#L392-L406), so disable_serialization() always runs while dispatches are still being instrumented. Check this ordering against the other services that hang off the same block.
No GPU drain in ATT stop; only disable_serialization (core.cpp#L609-L620).
collects_on is checked before pre_kernel_call (queue_hooks.cpp#L62-L65).
Refcounted serialization: does the all + per_agent split behave for a mix of scoped and unscoped callers? (queue_controller.hpp#L155-L183)
Disjoint ATT contexts: per_agent_scoping_test.cpp — intersects, disjoint start, serialization scope, write_hook agent filter.
Hook unit tests cover the gate and the completion-routing contract: queue_hooks_test.cpp — see StopContextInFlightCompletionRoutesViaHookPath and CompletionRoutingStaysWithTheProducingTracer.

Questions to Answer:

Does ATT-only profiling still intercept dispatches?
Does a GPU-1-only context leave GPU-0 unserialized?
Can two disjoint ATT contexts start together?
Is the enter-active / exit-registered split the right contract, or should exit also be scoped?

Out of scope: device ATT hook migration; deleting Queue::_callbacks; counters/SPM/PC-sampling migrations.

Issue Tracking

AIPROFSDK-1016

Test Plan

Committed coverage, all under source/lib/rocprofiler-sdk/thread_trace/tests/:

queue_hooks_test.cpp
- IsAnyActiveReturnsFalseWhenNoContextActive — the interceptor gate is off when no ATT context is active.
- StopContextInFlightCompletionRoutesViaHookPath — a dispatch tagged before stop_context is still collected afterwards, which is the property the registered-contexts exit path exists to provide.
- CompletionRoutingStaysWithTheProducingTracer — a packet is only consumed by the tracer that produced it, so overlapping tracers cannot cross-talk.
per_agent_scoping_test.cpp
- collects_on_and_intersects_match_configured_agents
- disjoint_agent_contexts_can_be_active_together
- overlapping_agent_contexts_still_conflict
- serialization_is_scoped_to_the_contexts_agents — asserts a context scoped to one agent does not serialize the others.
- write_hook_ignores_dispatches_on_other_agents
The existing thread-trace-api-* and rocprofv3-test-att-* integration suites exercise the positive hook path end to end.

Test Result

Treat the Checks tab as source of truth. The behavioural claims above are each asserted by a
named committed test rather than by a manual run.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

Related PRs

Earlier versions of this change:

#8790: the fork-hosted PR this change was moved from; its review threads still apply. It shows as merged because its base was this PR's head branch, not develop.
#11831: a conflict-resolved version of refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry #8790 against develop; closed shortly after opening.
#11964: an earlier rehost of refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry #8790 into ROCm/rocm-systems; closed as superseded and recreated as this PR from the --with-cursor-fixes branch.
#8586: the original all-services draft of the callback removal; this PR carries its thread-trace slice.
#8730: draft that resolves conflicts in Remove Callbacks from rocProfiler-SDK #8586 and modernizes it; this PR carries its thread-trace slice.

Parts of this change that landed as separate PRs on the earlier users/vkale/remove-callbacks-thread-trace branch:

#9934: scopes dispatch thread trace to configured agents; merged.
#11867: moves the queue hook packet tags off the callback registry id range; merged.

Other PRs in the remove-callbacks series: #11968 SPM, #11969 PC sampling, #11970 counters. Combined view: #12135.

therock-pr-bot · 2026-09-21T20:04:28Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot

Copilot review overview

🟡 Changes recommended

Critical synchronization, completion-routing, lifecycle, and test-compilation issues remain unresolved.

Get a fresh assessment by requesting another Copilot review.

Review effort: Lite
Findings: 5 · 2 · 1

Open (8)

ATT overlap check races with context publication · New
User callback executes while holding agents map lock · New
Tracer mutex acquisition can deadlock · New
Stopped contexts can consume later completions · New
Test lacks declaration for construct_agent_cache · New
Serializer state update races with refcount transitions · New
Context stop drains unrelated HSA queues · New
Completion routing declaration omits registered contexts · New

What changed in this PR

This PR migrates dispatch thread-trace profiling from per-queue callbacks to explicit hooks with agent-scoped serialization.

Changes:

Adds enqueue/completion hooks and packet provenance tags.
Implements per-agent serialization and overlap detection.
Adds tests, build wiring, and a CMake compatibility pin.

File	Reviewed changes and final findings
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/queue_hooks_test.cpp`	Tests hook gating. Critical (1): Missing `agent.hpp` include may prevent compilation.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/per_agent_scoping_test.cpp`	Tests agent scoping, conflicts, serialization, and filtering.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/CMakeLists.txt`	Registers test targets.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/queue_hooks.hpp`	Declares explicit hooks. Nit (2): Documentation does not reflect registered-context completion routing.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/queue_hooks.cpp`	Implements hook routing. Critical (2): Stopped contexts may claim later same-agent completions.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.hpp`	Adds agent-scoping and lifecycle APIs. Nit (1): Const methods mutate a non-mutable atomic.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.cpp`	Implements lifecycle and serialization. Critical (1): Lock ordering can deadlock; Critical (1): user callbacks run while holding the agent-map lock. Moderate (3): Stop synchronizes all queues instead of configured agents.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt`	Adds hook sources.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp`	Integrates explicit thread-trace hooks.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/CMakeLists.txt`	Registers hook headers.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp`	Defines packet producer tags.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp`	Declares scoped serialization APIs.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp`	Implements serialization refcounting. Moderate (2): Transition state can become stale before serializer updates. Moderate (1): Queue creation can race with disable.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt`	Includes queue-hook wiring.
`projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp`	Adds ATT conflict handling. Critical (1): Conflict checking races slot publication. Moderate (1): Teardown occurs before clearing the active slot.
`.github/workflows/rocprofiler-sdk-rocm_release_compatibility.yml`	Pins CMake for compatibility builds.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

…1967 Marker commit only; no file changes. It keeps the fork-side thread trace pull request open as a distinct step-1 PR stacked on users/vkale/remove-callbacks-thread-trace--with-cursor-fixes, which already contains every commit on this branch. This replaces ROCm#8790, which GitHub auto-closed as merged when its base branch came to contain its head commit. Merged pull requests cannot be reopened, so the step-1 PR is recreated here. Nothing reached develop. Co-authored-by: Cursor <cursoragent@cursor.com>

amd-vkale · 2026-09-21T20:34:20Z

Review items carried over from #8790

#8790 was the original thread trace migration PR. It was closed as merged by accident on 2026-09-21 when a compile fix was pushed to what had just become its base branch; it never merged into develop. Merged PRs cannot be reopened, so this PR is where the code is now reviewed, and #11973 is the replacement step-1 PR stacked on it. #11973 carries the full history; this comment is the actionable subset.

This PR needs fresh approval. @ApoKalipse-V, @MythreyaK and @bwelton approved #8790 back in July, but a substantial amount of work has landed since, so those approvals are stale.

Already addressed, for reviewers who saw the earlier discussion

Multi-context ATT ownership (raised by @bwelton). context.cpp now rejects a second dispatch_thread_trace context whose agent set intersects an active one, so overlapping contexts are refused while disjoint ones are allowed. This is stronger than the documented-invariant minimum that was asked for.
Dead client global in thread_trace/core.cpp is removed.
Hook tag range (client_ids.hpp). Now enum client_id : int64_t with negative values and a static_assert, keeping the tags disjoint from the positive ids QueueController::add_callback hands out. -1 is left unused because add_callback uses it as an "unassigned" initializer.
Build fix. DispatchThreadTracer::enabled is now mutable, because start_context() and stop_context() are const on the context interface. Without this the branch does not compile.

Still open

Serializer creation can start from a stale refcount snapshot. In serializer(), the refcount is read before the serializer lock is taken, so a concurrent enable_serialization can increment the refcount while the agent is still absent from _profiler_serializer — that update finds no entry to transition, and creation then inserts an entry built from the stale snapshot. The inverse can leave a serializer enabled after a disable.
Refcount update and serializer transition are applied under different locks. update_serialization() computes the transition list inside the refcount write lock, which is good, but applies the transitions after releasing it under a separate _profiler_serializer.wlock. An enable and a disable can therefore land their transitions in the opposite order from their refcount updates, leaving the serializer enabled with a final refcount of zero.

These two have a ready fix. refactor(rocprofiler-sdk): migrate SPM off the per-queue callback registry (moved PR 8887 to here) #11968 (SPM) already holds the refcount read lock across the entire lookup-and-insertion in serializer(), taking the locks in the same order update_serialization() does. Porting that here closes both. Worth deciding whether to port it onto this branch or to land refactor(rocprofiler-sdk): migrate SPM off the per-queue callback registry (moved PR 8887 to here) #11968 first and rebase.
Completion routing is not tied to the producing tracer. Every active ATT tracer sees every ATT completion, and post_kernel_call decrements post_move_data before establishing that the completing agent belongs to that tracer. The per-agent conflict check in item 1 of the addressed list narrows this a lot, since two contexts can no longer share an agent, but the hook still routes by service-wide client tag rather than by owning context. The open question is whether to persist a per-context owner with each injected packet, or to treat the admission-time check as the invariant and document it as such.
Unit coverage of the hook true path, originally raised by @ApoKalipse-V on refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry #8790. thread-trace-per-agent-test now covers collects_on and intersects against configured agents, disjoint versus overlapping contexts, serialization scoping, and write-hook filtering by agent. What is still missing is a mock or fake queue that exercises write_hook / signal_completion_hook with is_any_active() true; today that path is only reached through the rocprofv3 --att integration test, which needs a live HSA queue and dispatch. @bwelton, your view on whether that gap matters was invited on the original thread and is still open.

amd-vkale · 2026-09-21T21:20:21Z

@copilot Fix the code for all comments in this review thread.

When a review comment includes a suggested change, apply the suggestion exactly.

Do not make changes beyond what is described in the linked review thread.

Copilot · 2026-09-21T21:26:13Z

@copilot Fix the code for all comments in this review thread.

When a review comment includes a suggested change, apply ...

Addressed in 54c51c9.

amd-vkale · 2026-09-22T21:00:30Z

@copilot please reoslve all conflicts

Copilot

Copilot encountered an error and was unable to review this pull request. You can try again by re-requesting a review.

Copilot · 2026-09-22T21:04:15Z

@copilot please reoslve all conflicts

Merged develop into this branch in 38f2242. Git resolved it cleanly, so there were no manual conflict edits to make.

Refreshes the combined preview to develop f939bbe plus, in landing order, #11969 at 7f9d53c, #11968 at 9dd6af3, #11970 at 94925d9 and #11967 at 30bd96c. The tree is exactly the result of merging those heads onto develop; conflicts in hsa/queue.cpp, context/context.cpp and thread_trace/core.cpp were resolved additively. This picks up, since the last refresh: - the context startup marker on #11968 (cb5bf4e, 9dd6af3), now identical to 9d91611 on #11970 - #11967's merge of develop, including the thread trace resource_mode change (6fd8c2b), and its queue-hook test fix (30bd96c) Co-authored-by: Vivek Kale <vivek.kale@amd.com>

…back registry (#11969) Stop routing PC sampling kernel completions through the HSA queue controller's per-queue callback registry. Instead, call `pc_sampling::kernel_dispatch_phase_exit_hook` explicitly from `AsyncSignalHandler` in `hsa/queue.cpp`, which forwards to `kernel_completion_cb` after verifying the service is configured on the dispatch's agent. Add `pc_sampling::is_configured_on_agent` to the `no_real_consumers` gate in `WriteInterceptor`, replacing the notifier count that the old `add_callback` registration used to bump. This keeps marker-packet injection alive for PC-sampling-only runs on configured agents. Both hooks live in a new `queue_hooks.cpp`, added before the HSA-runtime-version early return in `pc_sampling/CMakeLists.txt` and guarded internally by `ROCPROFILER_SDK_HSA_PC_SAMPLING`, so they link as no-ops when HSA PC sampling support is unavailable. This is completion-only: marker injection stays inline in `WriteInterceptor`, with no batching or enqueue-hook change. Counters, thread trace, and SPM keep using the registry until their sibling PRs land. Hook names follow the series' `kernel_dispatch_phase_*_hook` convention. Other changes: - Move `kernel_completion_cb` out of the anonymous namespace and declare it in `hsa_adapter.hpp`; make its `kernel_pkt` param const. - Correct its comment: it runs on ROCr's async signal handler thread, not in the write interceptor. - Remove `intercept_cb_id` from `PCSAgentSession` (no longer needed). - Declare `is_hsa_initialized()` in `service.hpp` (was missing). - Remove stale `kernel_completion_cb` extern from test internals header. - Add unit tests for unconfigured gate, null-session no-op, and seeded-session configured path. Part of the per-service callback removal series. * Sibling PRs: #11967 (thread trace), #11968 (SPM), #11970 (counters) * Combined view: #12135 * Related: #8895, #8586, #8730 * Includes the completion-hook fix reviewed in #11871. JIRA ID: AIPROFSDK-1019 Co-authored-by: Cursor <cursoragent@cursor.com>

…istry (#11968) Stop registering dispatch SPM with the HSA queue controller's per-queue callback registry. The write interceptor and `AsyncSignalHandler` now call `spm::kernel_dispatch_phase_enter_hook` and `spm::kernel_dispatch_phase_exit_hook` explicitly. The enter hook walks active contexts and tags each SPM packet with `SPM_CLIENT_ID`; the exit hook walks registered contexts and picks out its own packets by that tag, so dispatches still in flight when a context stops complete normally. Add `spm::is_active_on_agent` to the `no_real_consumers` gate in `WriteInterceptor`, scoped to the queue's agent, replacing the notifier count the old registration used to bump, and keep per-packet mode while SPM is active on that agent. Scope SPM to GPU agents with the new experimental `rocprofiler_spm_dispatch_counting_service_set_agents()`. Two SPM contexts may be active together when their agent sets are disjoint. Queue-controller serialization is now reference counted per agent, so an agent stays serialized until every service that asked for it (counters, thread trace, SPM) has released it. SPM stop drains the GPU before releasing its hold. Thread trace's `resource_deinit` now releases serialization only if the tracer still held it, so it can no longer drop another service's hold. Other changes: - Read the context registry off a published snapshot, and serialize context start/stop transitions so lifecycle calls cannot act on a half-started or half-stopped context. - Take the contexts mutex in `spm::set_dispatch_agents`. - Add `hsa/queue_hooks/client_ids.hpp`: negative producer tags for migrated services, disjoint from the registry's positive ids. - Guard `rocprofiler_spm_destroy_counter_config` on the aqlprofile ABI. - Add tests: in-flight completion after stop, agent restriction, disjoint and overlapping contexts, concurrent start/stop, the write-interceptor gate, and thread trace's foreign-owner release. Counter collection (#11970) and thread trace (#11967) land after this; PC sampling migrated in #11969. Moved here from the fork-hosted #8887, which keeps the review history. Related: #8586, #8730, #7960 (kernel replay, AIPROFSDK-68) Sibling PRs: #11967 (thread trace), #11970 (counters) Combined view: #12135 JIRA ID: AIPROFSDK-1018 Co-authored-by: Cursor <cursoragent@cursor.com>

…--with-cursor-fixes # Conflicts: # projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp # projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp # projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.cpp Co-authored-by: Vivek Kale <vivek.kale@amd.com>

…o the kernel replay integration # Conflicts: # projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp # projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp Co-authored-by: Vivek Kale <vivek.kale@amd.com>

hsa/queue.cpp: thread trace (#11967) and counter collection (#11970) both call their hooks explicitly at the same places; keep both, in the order the vision branch uses (PC sampling, SPM, counters, thread trace). Co-authored-by: Cursor <cursoragent@cursor.com>

develop now contains SPM (#11968, e327f70) and PC sampling (#11969, 369f62d), so the preview is rebuilt as develop 3a6a0d8 plus the two services still in review: - thread trace #11967 at 369faf3 (merged cleanly); - counter collection #11970 at 9517f55 with develop merged in (64b1c4d), which resolves #11970's conflicts with the SPM and PC sampling changes that landed: in hsa/queue.cpp the counters hooks join the PC sampling and SPM hooks at the same call sites, and in context/context.cpp the counters and SPM intersects() checks coexist. hsa/queue.cpp between #11967 and #11970 resolves the same way, in the order this branch already used (PC sampling, SPM, counters, thread trace). The tree is exactly that integration (3b76ceb). This commit only records the previous preview head as a second parent, so the branch moves forward without a force push. Co-authored-by: Cursor <cursoragent@cursor.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>

Brings in "destroy thread trace agents off the async signal handler thread": 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 any more, so the merge also drops the parentheticals at the SPM, counter collection and thread trace enter hooks in hsa/queue.cpp that said other services still flow through signal_callback. Co-authored-by: Cursor <cursoragent@cursor.com>

The preview is develop 3a6a0d8 plus the two services still in review, as before, with thread trace #11967 now at f5eb45a: thread trace agents are no longer destroyed on ROCr's async signal handler thread, where on the HSA backend the process could hang at exit. With all four services migrated nothing calls QueueController::add_callback, so the enter-hook comments in hsa/queue.cpp no longer say other services flow through signal_callback. Validated on Banff (2 MI325X per run): full suite in CI and beta mode and 5 service-subset repetitions; every failure not shared with develop's runs was rerun 3x and is a known intermittent family. 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>

…read trace branches Only for use once #11968 and #11967 carry the start-side marker: update the comparison table and the known gap that said the marker existed only on this branch. Co-authored-by: Cursor <cursoragent@cursor.com>

…hooks docs Adds, on top of the previous preview tree: - the PC sampling start-window fix prepared for develop (cursor/context-start-window-pcs-b2ea): PC sampling now starts before context::start_context releases its start-in-progress marker, so a concurrent stop can no longer leave sampling enabled on an inactive context; with its race test and the SPM set_dispatch_agents race test; - the queue-hooks design pages from the #11970 and #11967 docs branches, revised for the end state: all four services on hooks, nothing registering through QueueController::add_callback, PC sampling under the marker. Validated on Banff (GPU pairs without the faulty GPU 3): full suite in CI and beta mode, every failure also seen on develop or in a known intermittent family; the four lifecycle race tests 20/20; docs render with Sphinx, the four pages in the nav. Co-authored-by: Cursor <cursoragent@cursor.com>

…--with-cursor-fixes

…th-cursor-fixes Brings in #11967 (thread trace off the per-queue callback registry). The conflicts in context/context.cpp, hsa/queue.cpp and thread_trace/core.cpp are where both PRs edit the same lines: thread trace keeps develop's code, and counter collection's hooks, fast-path gate, per-packet mode and agent-scoped conflict check sit next to it, as SPM's and PC sampling's already do. context.cpp also keeps develop's device SPM stop (#11879). Co-authored-by: Vivek Kale <vivek.kale@amd.com>

…a-hooks--with-cursor-help #11967 and #11970 have landed on develop as squash commits, while this branch carried their original commits, so the files they added conflict as add/add. In every conflict the two sides differ only by this branch's toggle change, so the branch's side is kept: the result is develop plus dfe73ff's change. Co-authored-by: Vivek Kale <vivek.kale@amd.com>

systems-assistant · 2026-10-09T23:38:44Z

TheRock Submodule Bump Activity

Newest first

2026-10-09 — Included in TheRock via ROCm/TheRock#8885.

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>

…lback registry (#11967) Stop registering dispatch thread trace (ATT) with the HSA queue controller's per-queue callback registry. The write interceptor and `AsyncSignalHandler` now call `thread_trace::kernel_dispatch_phase_enter_hook` and `thread_trace::kernel_dispatch_phase_exit_hook` explicitly. Add `thread_trace::is_active_on_agent` to the `no_real_consumers` gate in `WriteInterceptor`, scoped to the dispatch's agent, replacing the notifier count the old registration used to bump, and keep per-packet mode while ATT is active on that agent. An ATT context configured for one GPU no longer pulls queues on other GPUs off the fast path. Scope dispatch ATT to its configured GPU agents: serialization is enabled and released for those agents only (reference counted per agent), and two dispatch ATT contexts may be active together when their agent sets are disjoint; overlapping ones still conflict. Make completions safe against teardown: - `resource_deinit()` destroys the per-agent tracers only after releasing the agent-map lock, never on ROCr's async signal handler thread: an agent's destructor calls hsa_queue_destroy(), which waits on that thread, so destroying it there hung the process at exit. - A completion that arrives after `resource_deinit()` drops its trace but still drains the pending-completion count. - Keep SERIALIZE_ALL serialization barriers for dispatches that reach `pre_kernel_call()` while a context is stopping. - `stop_context()` and `resource_deinit()` release serialization only if the tracer still holds it. Add tests: per-agent scoping (disjoint and overlapping contexts, agent-scoped serialization) and queue-hook routing (in-flight completion after stop, a completion after teardown still drains, routing to the producing tracer, SERIALIZE_ALL while stopping). Counter collection still uses the registry until its sibling PR lands; PC sampling and SPM migrated in #11969 and #11968. Moved here from the fork-hosted #8790, which keeps the review history. Related: #8586, #8730 Sibling PR: #11970 (counters) Combined view: #12135 JIRA ID: AIPROFSDK-1016 Co-authored-by: Cursor <cursoragent@cursor.com>

…released Record already-merged remove-callbacks (#11967–#11970), SPM set_agents, and the #11335 kernel-iteration-range fix under Unreleased for ROCm 10.2, alongside the rocpd replay-pass aggregation note from this PR.

Resolve the CHANGELOG conflict by keeping the rocprofv3 --doctor (#12889) and hsa_amd_queue_create (#13720) entries from develop alongside this PR's Unreleased entries. Replace the queue-interposition, SPM agent-scoping, and kernel replay iteration-range entries with more detailed versions (#11967-#11970, #11335). Co-Authored-By: Claude <noreply@anthropic.com>

Copilot AI lite review requested due to automatic review settings September 21, 2026 20:03

amd-vkale requested review from a team as code owners September 21, 2026 20:03

github-actions Bot added project: rocprofiler-sdk github actions Pull requests that update GitHub Actions code labels Sep 21, 2026

Copilot started reviewing on behalf of amd-vkale September 21, 2026 20:04

View session

Copilot AI reviewed Sep 21, 2026

View reviewed changes

amd-vkale requested review from ApoKalipse-V, MythreyaK and bwelton September 21, 2026 20:12

systems-assistant Bot added the organization: ROCm label Sep 21, 2026

amd-vkale mentioned this pull request Sep 21, 2026

refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry (step 1) #11973

Closed

amd-vkale force-pushed the users/vkale/remove-callbacks-thread-trace--with-cursor-fixes branch from 0987bf0 to 58c06df Compare September 21, 2026 20:30

amd-vkale mentioned this pull request Sep 21, 2026

refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry #8790

Merged

1 task done

Copilot started work on behalf of amd-vkale September 21, 2026 21:21

View session

Copilot finished work on behalf of amd-vkale September 21, 2026 21:26

cursor Bot mentioned this pull request Sep 22, 2026

refactor(rocprofiler-sdk): migrate counter collection off the per-queue callback registry #11970

Merged

1 task done

amd-vkale requested a balanced review from Copilot September 22, 2026 21:01

Copilot AI reviewed Sep 22, 2026

Copilot started work on behalf of amd-vkale September 22, 2026 21:01

View session

Copilot started reviewing on behalf of amd-vkale September 22, 2026 21:02

View session

Copilot finished work on behalf of amd-vkale September 22, 2026 21:04

amd-vkale mentioned this pull request Sep 29, 2026

[DO NOT MERGE] Overnight stress: latest remove-callbacks PRs + kernel-replay stress test #12468

Closed

bwelton approved these changes Sep 30, 2026

View reviewed changes

amd-vkale changed the title ~~refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry (with cursor fixes)~~ refactor(rocprofiler-sdk): migrate thread trace off the per-queue callback registry Oct 6, 2026

amd-vkale added a commit that referenced this pull request Oct 7, 2026

merge: #11967 at 369faf3

44588bf

bgopesh approved these changes Oct 7, 2026

View reviewed changes

amd-vkale mentioned this pull request Oct 8, 2026

refactor(rocprofiler-sdk): resolve kernel-replay pass toggles in the dispatch hooks #13083

Draft

1 task

Merge branch 'develop' into users/vkale/remove-callbacks-thread-trace…

df9cce1

…--with-cursor-fixes

amd-vkale merged commit d19b095 into develop Oct 9, 2026
129 of 133 checks passed

amd-vkale deleted the users/vkale/remove-callbacks-thread-trace--with-cursor-fixes branch October 9, 2026 04:45

bgopesh mentioned this pull request Oct 14, 2026

docs(rocprofiler-sdk): add 1.4.2 CHANGELOG entries for ROCm 10.2.0 #13452

Merged

1 task done

Conversation

amd-vkale commented Sep 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Underlying Problem

Notes on Bigger Picture

Technical Details

Hooks and call sites

Per-agent serialization is reference counted

Reviewer Guide

Issue Tracking

Test Plan

Test Result

Submission Checklist

Related PRs

Uh oh!

therock-pr-bot Bot commented Sep 21, 2026

✅ All Policy Checks Passed

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-vkale commented Sep 21, 2026

Review items carried over from #8790

Already addressed, for reviewers who saw the earlier discussion

Still open

Uh oh!

amd-vkale commented Sep 21, 2026

Uh oh!

Copilot AI commented Sep 21, 2026

Uh oh!

amd-vkale commented Sep 22, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI commented Sep 22, 2026

Uh oh!

Uh oh!

systems-assistant Bot commented Oct 9, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

7 participants

amd-vkale commented Sep 21, 2026 •

edited

Loading

amd-vkale commented Sep 22, 2026 •

edited

Loading