Repository navigation
Conversation
Add three new subsystem hook free functions in parallel to the existing
Queue::_callbacks map indirection. Counters, thread_trace, and pc_sampling
each get a queue_hooks.{hpp,cpp} that wraps today's callback bodies.
No production callers yet: WriteInterceptor still iterates _callbacks in
this commit. C2 will pivot WriteInterceptor to call these hooks directly,
C3 deletes the now-dead callback registration plumbing.
Added:
- hsa/queue_hooks/client_ids.hpp
compile-time tag constants (replace runtime ClientID)
- counters/queue_hooks.{hpp,cpp}
write_hook, signal_completion_hook, is_any_active wrapping
existing queue_cb / completed_cb
- thread_trace/queue_hooks.{hpp,cpp}
same shape, wrapping DispatchThreadTracer pre_kernel_call /
post_kernel_call
- pc_sampling/queue_hooks.{hpp,cpp}
maybe_marker_packet returns std::optional<rocprofiler_packet>
because the marker is spliced into transformed_packets directly,
not via inst_pkt; signal_completion_hook; plus always-available
wrapper is_configured_on_agent that delegates to the
#if-guarded is_pc_sample_service_configured
- hsa/tests/fake_queue.hpp
shared FakeQueue test header, replacing two inline copies
- new unit tests in {counters,thread_trace,pc_sampling}/tests/
queue_hooks_test.cpp covering is_any_active / SymbolsExist /
Negative-paths; plus TEST(CountersQueueHooks, ...) and
TEST(CountersMultiCallback, ...) colocated in counters/tests/core.cpp
Modified:
- counters/tests/core.cpp, tests/hsa_barrier.cpp now include the
shared fake_queue.hpp instead of defining FakeQueue inline
- pc_sampling/hsa_adapter.{hpp,cpp} promoted kernel_completion_cb
from anonymous namespace + declared in header so
signal_completion_hook can call it
- per-subsystem CMakeLists.txt entries wire new files into the build
- pc_sampling/tests/CMakeLists.txt links counter-test-constants for
the shared FakeQueue dependency
Verified: 326 pass / 4 fail unittests on banff TheRock gfx942 container.
The 4 failures are pre-existing att_decoder_waitcnt_test {gfx9, gfx10,
gfx12, fail_conditions} unrelated to this refactor.
Replace WriteInterceptor's early-return predicate to use the new
is_any_active queries instead of queue.get_notifiers(), so the predicate
keeps working after callback installation is deleted in this same commit.
Replace Queue::_callbacks map iteration in WriteInterceptor and
AsyncSignalHandler with direct calls to the C1 hook scaffolding:
- counters::write_hook + counters::signal_completion_hook
- thread_trace::write_hook + thread_trace::signal_completion_hook
- pc_sampling::maybe_marker_packet + pc_sampling::signal_completion_hook
Replace the should_batch_packets map-AND-fold with the hardcoded
expression !counters::is_any_active() && !thread_trace::is_any_active().
Replace the #if ROCPROFILER_SDK_HSA_PC_SAMPLING block at the marker
emission site in queue.cpp with a call to pc_sampling::maybe_marker_packet.
Same splice position, same configured-vs-active semantics; the disabled-
build #if guard moves into the hook implementation.
Delete callback installation sites:
- counters/core.cpp: the for-loop installing per-counter_callback_info
queue_callbacks_t entries via controller->add_callback. The
enabled bool / dispatch_counter_collection mutations stay because
they drive the new is_any_active query (which filters on
dispatch_counter_collection != nullptr).
- thread_trace/core.cpp: DispatchThreadTracer::start_context and
stop_context bodies reduced to just enable/disable_serialization.
The client field stays (unused) until C3 deletes it.
- pc_sampling/hsa_adapter.cpp: the controller->add_callback call and
its lambda payload. Surrounding PC sampling setup is preserved.
The queue_callbacks_t map, ClientID typedef, and the add_callback /
register_callback / signal_callback machinery still exist but are no
longer populated or read. C3 deletes the dead infrastructure.
Update counters/tests/core.cpp: in TEST(core, start_stop_buffered_ctx)
and TEST(core, start_stop_callback_ctx), the iterate_callbacks-finds-
queue_id assertions are removed (replaced with TODO(C3) comments).
These assertions tested the installation mechanism rather than
behavior; the rest of each test (context state, dispatch_counter_
collection->callbacks config, enabled flag transitions) still runs and
still validates the intended behavior.
Add tests/should_batch_packets_test.cpp pinning the new should_batch_
packets contract:
- BatchesPacketsWhenNoSubsystemActive (idle path).
- RequiresPerPacketWhenCountersActive (counters path; uses the
real registered-context activation pattern from
start_stop_buffered_ctx).
- RequiresPerPacketWhenATTActive (SKIP'd with documented reason;
coverage delegated to the rocprofv3 --att C2 manual smoke).
Verified: 330 pass / 4 known fail unittests on banff TheRock gfx942
container. The 4 failures are the pre-existing att_decoder_waitcnt_test
{gfx9, gfx10, gfx12, fail_conditions} unrelated to this refactor.
Following the WriteInterceptor pivot to direct subsystem calls in C2,
the callback map infrastructure has no remaining users. Delete it.
Deleted from hsa/queue.{hpp,cpp}:
- queue_callbacks_t struct
- queue_batch_packets_callback_t, queue_write_callback_t,
queue_completion_callback_t function typedefs
- write_cb_t, completed_cb_t aliases
- callback_map_t typedef
- Queue::_callbacks field (and its rlock-during-fan-out machinery)
- Queue::_notifiers atomic + get_notifiers accessor
- Queue::register_callback, Queue::remove_callback declarations + impls
- Queue::signal_callback template method
Deleted from hsa/queue_controller.{hpp,cpp}:
- QueueController::_callback_cache field
- QueueController::add_callback, QueueController::remove_callback,
QueueController::iterate_callbacks declarations + impls
- agent_callback_tuple_t, callback_iterator_cb_t, client_id_map_t
helper typedefs
- default_agent sentinel struct (~60 LOC, no remaining callers after
add_callback/iterate_callbacks deletion)
- add_queue simplified to just the queue-map insertion (the
callback-cache walk is gone)
Deleted from hsa/rocprofiler_packet.hpp:
- using ClientID = int64_t typedef
- inst_pkt_t element type changed from {AQLPacket*, ClientID} to
{AQLPacket*, int64_t} directly
Deleted from counters/:
- counters::ClientID typedef (counters/dispatch_handlers.hpp)
- counters::inst_pkt_t element type updated to int64_t
- counter_callback_info::queue_id field (counters/core.hpp)
- TODO(C3) comments left in tests/core.cpp by C2 (now resolved)
- One stale static_cast<counters::ClientID> in tests/core.cpp
Deleted from pc_sampling/:
- intercept_cb_id field from pc_sampling/types.hpp
- Stale kernel_completion_cb forward declaration in
pc_sampling/tests/pc_sampling_internals.hpp (used the removed
ClientID type, had no test consumers, and the production
kernel_completion_cb now lives in hsa_adapter.hpp with a different
signature)
Verified: 330 pass / 4 known fail unittests on banff TheRock gfx942
container. The 4 failures are pre-existing att_decoder_waitcnt_test
{gfx9, gfx10, gfx12, fail_conditions} unrelated to this refactor.
Repo-wide grep confirms zero remaining references to the deleted
symbols (queue_callbacks_t, ClientID, callback_map_t, _callback_cache,
register_callback, remove_callback, signal_callback, iterate_callbacks,
_callbacks, _notifiers, intercept_cb_id) outside the new
hsa/queue_hooks/ infrastructure and the unaffected
projects/rocprofiler-sdk-attach/ shim.
Net diff: 11 files changed, ~30 insertions(+), ~265 deletions(-).
…olding Remove path-to-self header comments, historical archaeology comments explaining what was deleted in earlier commits, and over-narrated test section headers. Code now reads on its own merits; commit history holds the why-it-changed narrative. Net: -230 / +79 lines, no semantic changes.
Textual conflicts resolved (develop <- users/bewelton/remove-callback): - hsa/queue.cpp: keep both includes; blend WriteInterceptor gate = 8586 is_any_active model + develop HIP-graph carve-out, drop removed get_notifiers() - counters/core.cpp: drop add_callback loop, keep develop callback_thread_start(); drop remove_callback teardown in stop_context (implicit under hook model) - hsa/queue_controller.cpp: simplified add_queue, keep develop create_queue_state() - thread_trace/core.cpp: take 8586 deletion KNOWN NOT-YET-COMPILING (modernization TODO, tracked for Ben/Mythreya): - spm/core.cpp: SPM (added to develop after this branch) still uses add_callback/queue_callbacks_t/remove_callback -> must be ported to an spm::write_hook/signal_completion_hook/is_any_active and wired into queue.cpp - spm/tests/core.cpp: iterate_callbacks usage needs updating Co-authored-by: Cursor <cursoragent@cursor.com>
SPM was added to develop after Ben's branch and still used the removed
callback registry (add_callback/queue_callbacks_t/remove_callback). Port it to
the same free-function hook pattern as counters/thread_trace/pc_sampling:
- new spm/queue_hooks.{hpp,cpp}: free functions write_hook /
signal_completion_hook / is_any_active over active dispatch_spm contexts
- spm/core.cpp: start/stop_context drop the add_callback/remove_callback loops
- spm/core.hpp: remove now-unused queue_id (also drops the removed hsa::ClientID)
- hsa/queue_hooks/client_ids.hpp: add SPM_CLIENT_ID
- hsa/queue.cpp: call spm::write_hook + spm::signal_completion_hook, add SPM to
should_batch_packets, and fix the WriteInterceptor gate to detect real
consumers explicitly (counters/ATT/SPM/PC-sampling) since get_notifiers() and
the registry are gone (full_packet_instrumentation_context_filter only covers
kernel-dispatch tracing)
Co-authored-by: Cursor <cursoragent@cursor.com>
Drop iterate_callbacks / queue_id assertions (registry removed); rely on the context enabled flag for activeness checks. Co-authored-by: Cursor <cursoragent@cursor.com>
amd-vkale
marked this pull request as ready for review
July 16, 2026 22:39
1 task done
This was referenced Jul 20, 2026
1 task done
amd-vkale
added a commit
to amd-vkale/rocm-systems
that referenced
this pull request
Jul 20, 2026
…back registry Isolates the PC-sampling slice of the callback-removal effort (reference PRs ROCm#8730 / ROCm#8586) into a small, single-service change, per review guidance to land callback removal one service at a time. Independent of kernel replay. PC sampling no longer registers a queue_callbacks_t with the HSA queue controller in pc_sampling_service_finish_configuration. Instead: - pc_sampling::signal_completion_hook (new pc_sampling/queue_hooks.{hpp,cpp}) delivers kernel completion to the CID manager; it is called explicitly from the HSA async signal handler in hsa/queue.cpp. kernel_completion_cb is moved out of the anonymous namespace and declared in hsa_adapter.hpp. - pc_sampling::is_configured_on_agent replaces the queue.get_notifiers() signal the old registration used to provide, so the write interceptor still injects the marker packet for PC-only runs. It is always linkable (guarded by ROCPROFILER_SDK_HSA_PC_SAMPLING) and compiled before the pc_sampling CMake early-return so it links even without HSA PC sampling support. The marker packet injection in the write interceptor is unchanged (it was never registry-based). Removes the now-unused intercept_cb_id from PCSAgentSession and adds a queue_hooks unit test. The per-queue callback registry is retained for the remaining service (thread-trace/spm/counters migrate in their own PRs). Co-authored-by: Cursor <cursoragent@cursor.com>
amd-vkale
added a commit
that referenced
this pull request
Aug 13, 2026
…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>
1 task
cursor Bot
pushed a commit
that referenced
this pull request
Aug 14, 2026
…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>
cursor Bot
pushed a commit
that referenced
this pull request
Aug 14, 2026
…istry Isolates the SPM slice of the callback-removal effort (reference PRs #8730 / removal one service at a time. Independent of kernel replay. SPM 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: - spm::write_hook / signal_completion_hook / is_any_active (new spm/queue_hooks.{hpp,cpp}) iterating active dispatch_spm contexts and calling the existing pre_kernel_call / post_kernel_call - hsa/queue_hooks/client_ids.hpp: stable inst_pkt producer tags The per-queue callback registry is intentionally retained for counters, thread-trace and pc-sampling; those migrate in follow-up PRs. The WriteInterceptor gate and should_batch predicate detect SPM via is_any_active() (it no longer contributes to queue.get_notifiers()). Removes the now-unused queue_id from spm_counter_callback_info and updates spm tests off iterate_callbacks. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
|
This pull request has been automatically closed due to inactivity (30 days with no updates). If you'd like to continue working on this, feel free to reopen the PR or create a new one. |
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
This was referenced Sep 21, 2026
Merged
Contributor
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
amd-vkale
added a commit
that referenced
this pull request
Oct 2, 2026
…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>
Contributor
|
This pull request has been automatically closed due to inactivity (30 days with no updates). If you'd like to continue working on this, feel free to reopen the PR or create a new one. |
amd-vkale
added a commit
that referenced
this pull request
Oct 5, 2026
…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>
amd-vkale
added a commit
that referenced
this pull request
Oct 9, 2026
…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>
amd-vkale
added a commit
that referenced
this pull request
Oct 9, 2026
…ue callback registry (#11970) 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>
anujshuk-amd
pushed a commit
that referenced
this pull request
Oct 12, 2026
…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>
anujshuk-amd
pushed a commit
that referenced
this pull request
Oct 12, 2026
…ue callback registry (#11970) 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>
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
Fixes conflicts and modernizes PR #8586
Technical Details
Issue Tracking
Test Plan
Test Result
Submission Checklist