[DO NOT MERGE] Remove callbacks from rocProfiler-SDK: integration preview of all four services by amd-vkale · Pull Request #12135 · ROCm/rocm-systems

amd-vkale · 2026-09-23T18:02:42Z

Integration preview only, not for merge. develop plus the four service migrations merged together to show the combined end state. Review comments belong on PRs 11967, 11968, 11969 and 11970.

Issue Tracking

JIRA ID : AIPROFSDK-68

therock-pr-bot · 2026-09-23T18:03:35Z

✅ 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?

therock-pr-bot · 2026-09-23T18:03:37Z

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

❌ PR Description

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

amd-vkale · 2026-09-23T18:03:54Z

@copilot please resolve all conflicts

amd-vkale · 2026-09-23T18:04:20Z

@copilot resolve the merge conflicts in this pull request

Copilot · 2026-09-23T18:09:40Z

@copilot please resolve all conflicts

Resolved in 23a471d.

Copilot

Copilot review overview

🟡 Changes recommended

Concurrent activation and completion-routing defects can create conflicting contexts and misroute thread-trace packets.

Get a fresh assessment by requesting another Copilot review.

Review effort: Balanced
Findings: 2 · 1 · 1

Open (4)

Context conflict check and publication are not atomic · New
Active agent mutation races with context activation · New
Packet routing can misattribute data across tracer contexts · New
Obsolete known gap text contradicts implemented hooks · New

What changed in this PR

Integration preview combining callback-registry removal for counters, SPM, thread trace, and PC sampling.

Changes:

Replaces per-queue callbacks with explicit dispatch hooks.
Adds per-agent collection, conflict, and serialization handling.
Adds regression, integration, and performance tests.

File	Description
`tests/queue-hooks-perf/run_launch_sweep.py`	Runs launch-count performance sweeps.
`tests/queue-hooks-perf/run_and_validate.py`	Validates performance bounds.
`tests/queue-hooks-perf/perf_cost_model.py`	Defines performance ceilings.
`tests/queue-hooks-perf/main.cpp`	Adds the HIP performance workload.
`tests/queue-hooks-perf/CMakeLists.txt`	Registers performance tests.
`tests/counter-collection-per-agent/per_agent_scoping_app.cpp`	Adds multi-GPU scoping integration coverage.
`tests/counter-collection-per-agent/CMakeLists.txt`	Builds the scoping integration test.
`tests/CMakeLists.txt`	Enables the new test directories.
`source/lib/rocprofiler-sdk/thread_trace/tests/queue_hooks_test.cpp`	Tests thread-trace completion routing.
`source/lib/rocprofiler-sdk/thread_trace/tests/per_agent_scoping_test.cpp`	Tests thread-trace agent scoping.
`source/lib/rocprofiler-sdk/thread_trace/tests/CMakeLists.txt`	Registers thread-trace tests.
`source/lib/rocprofiler-sdk/thread_trace/queue_hooks.hpp`	Declares thread-trace hooks.
`source/lib/rocprofiler-sdk/thread_trace/queue_hooks.cpp`	Implements thread-trace routing.
`source/lib/rocprofiler-sdk/thread_trace/core.hpp`	Adds agent-set helpers.
`source/lib/rocprofiler-sdk/thread_trace/core.cpp`	Migrates thread trace from callbacks.
`source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt`	Builds thread-trace hooks.
`source/lib/rocprofiler-sdk/spm/tests/queue_hooks_test.cpp`	Tests SPM hook filtering.
`source/lib/rocprofiler-sdk/spm/tests/core.cpp`	Expands SPM lifecycle and scoping tests.
`source/lib/rocprofiler-sdk/spm/tests/CMakeLists.txt`	Registers SPM hook tests.
`source/lib/rocprofiler-sdk/spm/service.cpp`	Exposes SPM agent scoping.
`source/lib/rocprofiler-sdk/spm/queue_hooks.hpp`	Declares SPM hooks.
`source/lib/rocprofiler-sdk/spm/queue_hooks.cpp`	Implements SPM hook routing.
`source/lib/rocprofiler-sdk/spm/core.hpp`	Declares SPM agent configuration.
`source/lib/rocprofiler-sdk/spm/core.cpp`	Adds scoped SPM lifecycle handling.
`source/lib/rocprofiler-sdk/spm/CMakeLists.txt`	Builds SPM hooks.
`source/lib/rocprofiler-sdk/pc_sampling/types.hpp`	Removes callback registration state.
`source/lib/rocprofiler-sdk/pc_sampling/tests/queue_hooks_test.cpp`	Tests PC-sampling hook behavior.
`source/lib/rocprofiler-sdk/pc_sampling/tests/CMakeLists.txt`	Registers PC-sampling hook tests.
`source/lib/rocprofiler-sdk/pc_sampling/queue_hooks.hpp`	Declares PC-sampling hooks.
`source/lib/rocprofiler-sdk/pc_sampling/queue_hooks.cpp`	Implements PC-sampling completion routing.
`source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.hpp`	Exposes kernel completion handling.
`source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp`	Removes callback registration.
`source/lib/rocprofiler-sdk/pc_sampling/CMakeLists.txt`	Builds always-linkable hooks.
`source/lib/rocprofiler-sdk/hsa/queue.cpp`	Invokes migrated services directly.
`source/lib/rocprofiler-sdk/hsa/queue_hooks/CMakeLists.txt`	Registers shared hook tags.
`source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp`	Defines service packet tags.
`source/lib/rocprofiler-sdk/hsa/queue_controller.hpp`	Adds scoped serialization state.
`source/lib/rocprofiler-sdk/hsa/queue_controller.cpp`	Implements serialization reference counting.
`source/lib/rocprofiler-sdk/hsa/CMakeLists.txt`	Includes queue-hook infrastructure.
`source/lib/rocprofiler-sdk/dispatch_counting_service.cpp`	Exposes counter agent scoping.
`source/lib/rocprofiler-sdk/counters/tests/queue_hooks_test.cpp`	Tests counter completion routing.
`source/lib/rocprofiler-sdk/counters/tests/per_agent_scoping_test.cpp`	Tests counter agent scoping.
`source/lib/rocprofiler-sdk/counters/tests/core.cpp`	Updates lifecycle assertions.
`source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt`	Registers new counter tests.
`source/lib/rocprofiler-sdk/counters/sample_processing.cpp`	Reference-counts the consumer thread.
`source/lib/rocprofiler-sdk/counters/queue_hooks.hpp`	Declares counter hooks.
`source/lib/rocprofiler-sdk/counters/queue_hooks.cpp`	Implements counter hook routing.
`source/lib/rocprofiler-sdk/counters/core.hpp`	Declares counter agent configuration.
`source/lib/rocprofiler-sdk/counters/core.cpp`	Migrates counters and adds scoping.
`source/lib/rocprofiler-sdk/counters/CMakeLists.txt`	Builds counter hooks.
`source/lib/rocprofiler-sdk/context/context.hpp`	Stores service agent sets.
`source/lib/rocprofiler-sdk/context/context.cpp`	Adds agent-aware context conflicts.
`source/include/rocprofiler-sdk/experimental/spm.h`	Documents the SPM scoping API.
`source/include/rocprofiler-sdk/dispatch_counting_service.h`	Documents the counter scoping API.
`source/docs/conceptual/queue_hooks/queue_callback_removal_architecture.md`	Documents the hook architecture.

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

cursor · 2026-09-25T02:22:24Z

Posting note: drafted with Cursor, but I (@amd-vkale) own this post.

Agreed, this was a real race. It's fixed on the counter PR (#11970) in 9d91611990b, and the preview picks it up as of c73bbff612f.

The conflict scan now runs under get_contexts_mutex(), after waiting out any in-flight start or stop. The new context is then marked pending until its slot is published and its services have started, so a concurrent overlapping start can't scan past it: context.cpp L528-L616.

cursor · 2026-09-25T02:22:24Z

Posting note: drafted with Cursor, but I (@amd-vkale) own this post.

Agreed. The SPM path lacked the guard the counter path has. It's fixed on the SPM PR (#11968) in a308ba361d3, and the preview picks it up as of c73bbff612f.

spm::set_dispatch_agents now takes get_contexts_mutex() and waits out in-flight starts and stops before touching the active state, matching counters: spm/core.cpp L342-L343. This can't deadlock against the context lookup, because that lookup reads a lock-free snapshot.

cursor · 2026-09-25T02:22:24Z

Posting note: drafted with Cursor, but I (@amd-vkale) own this post.

Agreed. Filtering by client id alone didn't keep provenance. It's fixed on the thread trace PR (#11967) in 33b1c5c1bca, and the preview picks it up as of c73bbff612f.

Each tracer gets a tracer_id that is never reused (core.hpp L186-L193). Every packet it builds is stamped with that id and holds a reference to the producing agent state. post_kernel_call skips packets it doesn't own and delivers through the packet's own reference, so it no longer looks anything up in the tracer's agent map (core.cpp L540-L565). ThreadTraceQueueHooks.CompletionRoutingStaysWithTheProducingTracer covers this: queue_hooks_test.cpp L195.

cursor · 2026-09-25T02:22:24Z

Posting note: drafted with Cursor, but I (@amd-vkale) own this post.

Agreed. The text was written from the counter PR's point of view and read as wrong in the combined preview. It's reworded on #11970 in 94925d9432a, and the preview picks it up as of c73bbff612f.

The gap is now framed as a condition: until SPM, thread trace and PC sampling land, each still registers through add_callback. Once all four are in, nothing calls it and the registry can go as a follow-up. The intro also names all four PRs and notes that they can land in any order: architecture doc L10-L11 and L207-L211.

Copilot · 2026-09-23T18:12:56Z

@copilot resolve the merge conflicts in this pull request

Conflicts are already resolved in 23a471d.

When multiple stop reason bits are set, the loop extracts them one by one into 'one_bit', but the switch statement was incorrectly using 'stop_reason' (the full bitmask) instead of 'one_bit' (the single bit). This caused multi-bit stop reasons to fall through to the default case and return 0 instead of the combined exception flags. Ticket: AIROCGDB-635

…eadability-redundant-member-init` (#12345) ## Motivation - Fixing clang-tidy issues and violations ## Technical Details - Run script for fixing `readability-redundant-member-init` ## Issue Tracking JIRA ID: AIPROFSYST-496 ## Test Plan - Existing tests pass. ## Test Result - Passing. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

…odernize-use-designated-initializers` (#12357) ## Motivation - Fixing clang-tidy issues and violations ## Technical Details - Run script for fixing `modernize-use-designated-initializers` ## Issue Tracking JIRA ID: AIPROFSYST-496 ## Test Plan - Existing tests pass. ## Test Result - Passing. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Signed-off-by: David Galiffi <David.Galiffi@amd.com> Co-authored-by: David Galiffi <David.Galiffi@amd.com>

…#12524) ## Motivation  coverage script is failing when we remove some files. ## Technical Details  fix script to skip missing files. ## Issue Tracking    JIRA ID: AIPROFSYST-785 ## Test Plan  Existing tests pass ## Test Result  All tests are passing. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

amd-argus · 2026-10-07T00:31:51Z

@cppccppccppc @sunguobao PTAL

…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>

…b2ea, 0a77ce8) Brings in the change prepared against develop: PC sampling now starts while its context is still marked as starting, so a concurrent stop can no longer claim the context before sampling is enabled and leave it enabled on a context that is not active. Also brings the two new race tests (pc_sampling.concurrent_start_stop_does_not_leave_sampling_enabled, spm_core.set_agents_does_not_race_context_start). Co-authored-by: Cursor <cursoragent@cursor.com>

…mpare the services List the architecture page in the conceptual section of the docs table of contents, with a glob entry for conceptual/queue_hooks/*_queue_hooks so each sibling PR can add its service page without editing the toc. Replace the claim that the other services mirror counter collection with a table of where SPM, thread trace and PC sampling actually differ, and state the drain bound and the start-marker gap precisely. 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>

The page was written while all four service migrations were open and compared their pull request heads. Counter collection, SPM and PC sampling now use explicit hooks and only thread trace still registers with the callback registry, so the page compares the services and states each behaviour as the code implements it. This corrects the thread trace column of the comparison, SPM's serialization and drain behaviour, the stop phases and the start-in-progress marker, which covers neither device counter collection nor PC sampling, the kernel replay section and the known gaps. Co-authored-by: Cursor <cursoragent@cursor.com>

Add a page beneath the callback registry removal page that describes how dispatch SPM is wired into dispatch interception: its hooks and call sites, which context set each phase reads, how a completion finds the callback that produced its packet, the start and stop ordering, per-agent gating and serialization, the unit tests and the known gaps. The page starts from the one written for the SPM pull request and is rewritten to match the code, in which a stop drains and releases serialization only when it clears the enabled flag and the race tests repeat 1000 times within each test. Co-authored-by: Cursor <cursoragent@cursor.com>

Add a page beneath the callback registry removal page that describes how PC sampling is wired into dispatch interception without the registry: the configuration gate, completion delivery through the agent's session, where start and stop run in the context lifecycle, the tests and the known gaps. The page starts from the one written for the PC sampling pull request and is rewritten for the current context lifecycle. PC sampling starts after the start-in-progress marker is released, so a concurrent stop can claim the context before sampling has started; the page describes what that leaves enabled. Co-authored-by: Cursor <cursoragent@cursor.com>

Add conceptual/queue_hooks/thread_trace_queue_hooks.md. It describes the dispatch thread trace enter and exit hooks and where hsa/queue.cpp calls them, which contexts each phase reads, how a completion is matched to the tracer that injected its packet, what start, stop and finalization guarantee, per-agent gating, the tests that cover each property, and the known gaps. The page is marked orphan and is not yet in any table of contents. Co-authored-by: Cursor <cursoragent@cursor.com>

…lict check context::start_context said overlapping thread trace agent sets would cross-talk in post_kernel_call. Completions are matched by the tracer id stamped on each packet, so they no longer can; keep the rule and say why a thread trace context, unlike an SPM one, never claims every agent. Co-authored-by: Cursor <cursoragent@cursor.com>

…s migrated The pages were written per pull request, each against a tree in which some service still used the per-queue callback registry. With thread trace, counter collection, SPM and PC sampling all on hooks, and PC sampling now starting under the start-in-progress marker: - the architecture page says nothing registers through QueueController::add_callback, gives thread trace's column in the service comparison (same hooks, registered-context exit hook, owner found by tracer id, agent-scoped serialization, conflict rule), and lists the registry itself as dead code to remove; - the start-marker description and the known gaps cover PC sampling starting before the marker is released; device counter collection is the one service left after it; - the PC sampling and SPM pages list the two new race tests, and the thread trace page drops its orphan front matter now that the toc glob includes it. 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>

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.

…p drain The architecture and SPM queue-hooks pages described hsa::queue_controller_sync() as the per-queue Queue::sync() slices only. It first waits, with no time limit, for the queue-interposition completion monitor's in-flight batches (interposition_sync()), so only the per-queue part of the drain is bounded. Same correction as the code comments in the previous commit.

mangupta

🤖 AI-assisted pre-review (for @mangupta)

[IMPORTANT] Kernel-node Set/GetAttribute now return hipErrorInvalidValue for every attribute other than four, with no changelog or doxygen update

Before this change, GraphKernelNode::SetAttrParams accepted any attribute ID and recorded it in kernelAttrInUse_. GetAttrParams returned hipSuccess for any attribute matching the in-use one. The new code returns hipErrorInvalidValue for every attribute except AccessPolicyWindow, Cooperative, Priority and ClusterDimension. Applications or frameworks that set other launch attributes on graph kernel nodes previously succeeded and now fail. This affects hipGraphKernelNodeSetAttribute, hipGraphKernelNodeGetAttribute and paths that copy attributes. Examples include memSyncDomain-style attributes that are mapped in hipLaunchAttribute but not stored here.

Suggested fix: Confirm the stricter behavior matches CUDA for each hipKernelNodeAttrID / hipLaunchAttributeID value. Where CUDA accepts an attribute (even as a no-op), keep accepting it. Otherwise, document the behavior change in CHANGELOG.md and the doxygen for hipGraphKernelNodeSetAttribute/GetAttribute, and add a negative test that pins which IDs are rejected.

[BLOCKING] Restore cross-rank agreement before selecting symmetric kernels

The scheduler no longer rejects partially registered buffers or reaches cross-rank consensus on symmetric-kernel eligibility, so ranks can select incompatible collective implementations.

Suggested fix: Restore the both-windows-registered check and a collective eligibility vote before finalizing symmetric tasks, or move an equivalent all-rank invariant check earlier in selection and prove every rank consumes the same decision.

[IMPORTANT] Size each shared completion queue for all attached QPs

Each physical rail creates one completion queue with capacity for only one QP's send and receive depths, but every QP assigned to that rail shares the queue. Configuring more than one QP per adapter can therefore overflow the CQ under legal in-flight load.

Suggested fix: Compute the required CQ depth per rail from the number of QPs attached to that rail and the maximum CQ-generating operations per QP, validate that aggregate against MaxCompletionQueueDepth, and either allocate that depth or reduce/reject the configured QP count.

[IMPORTANT] Retain GET source registrations until response transmission completes

A remote GET response queues a task containing pointers into a local memory registration without retaining that registration. The registration can be deregistered and freed before the queued payload is copied, causing a use-after-free or access through an unmapped GDRCopy handle.

Suggested fix: Acquire a task-owned reference on the source MR before enqueueing a successful GET response, release it when the send task completes, and also release it on every queue-abort, connection-error, and teardown path. Ensure deregistration cannot remove the handle while such a reference exists.

[IMPORTANT] Route per-rail CQ errors through the failover path

A completion-queue error on one NetworkDirect rail immediately makes the entire communicator fatal, even when failover is enabled and another rail remains healthy.

Suggested fix: For a CQ notification associated with devIndex, invoke ncclNdHandleRailFailure when failover is available, flush and retire that rail, and allow outstanding eligible requests to replay. Escalate to ncclNdFailDataPath only when failover is disabled, exhausted, or no operational rail remains.

[IMPORTANT] Link the SDMA factory test when only installed rocSHMEM is available

The new rccl-UnitTestsGinAnvilSdmaFactory target includes the required anvil.cpp implementation only when ROCSHMEM_SOURCE_DIR is defined. Configurations using an installed rocSHMEM through ROCSHMEM_INCLUDE_DIR or ROCSHMEM_INSTALL_DIR still create this target but provide only headers and low-level system libraries, leaving the Anvil implementation unavailable at link time.

Suggested fix: When ROCSHMEM_SOURCE_DIR is unavailable, link an installed rocSHMEM/Anvil target that exports the required implementation, locate an explicitly supplied source path, or disable the target with a clear configuration message. Keep target creation conditional on having both headers and implementation.

[IMPORTANT] Add the required SPDX identifiers to new files

Multiple new files do not carry the repository-required SPDX license header, so the pull request fails the supplied license checklist.

Suggested fix: Apply the repository-prescribed SPDX header to every file listed by the missing_spdx_new_files checklist, preserving any third-party attribution and license terms where vendored Microsoft headers require them, then rerun the license check.

[IMPORTANT] New mi455 configs reference a CheckMode test this PR renames elsewhere, so the entry now silently reports SKIPPED

This PR adds a symmetric_abort_checkmode_tests block to mi455_ainic_roce.json and mi455_ainic_roce_func.json. That block uses the filter SymCheckMode_Registration.Default_RegistrationMismatch_Accepted. The same PR replaces that entry in mi300x_mellanox_ib.json and mi355x_thor2_roce.json with SymCheckMode_Registration.Default_MatchingRegistration_Succeeds, which shows the old test was renamed.

Suggested fix: Update both mi455 configs to SymCheckMode_Registration_DefaultMatching / SymCheckMode_Registration.Default_MatchingRegistration_Succeeds, matching mi300x and mi355x.

[IMPORTANT] IB-CAST multithread suite is enabled on mi455, while every other NET IB suite there is disabled for lack of verbs userspace

The mi455 configs now disable all NET IB suites, including cast_wrr_split, cast_stress and net_ib_multithread_independent, with the reason "Broadcom Thor Ultra has no matching OFED/verbs userspace". They also set NCCL_IB_DISABLE=1 for the Ethernet transport. In the same change they add and enable "NET IB - CAST Multithread (4)" (config cast_multithread, which extends cast_base). That suite drives NetIbMPITest.CastStressMultiRoundTwoConns over IB-CAST.

Suggested fix: Set cast_multithread to "enabled": false in the mi455 configs, with the same Thor Ultra rationale, or explain why IB-CAST multithreading works where the other CAST suites do not.

[IMPORTANT] Dropping ASAN_OPTIONS=detect_leaks=0 for the Python launcher will make pytest ctest entries fail under TheRock ASAN

sanitizer_runtime_env lost its suppress_leaks parameter, so enable_sanitizer_python_launcher no longer injects ASAN_OPTIONS=detect_leaks=0. Under THEROCK_SANITIZER=ASAN|HOST_ASAN, the pytest-based ctest entries will report CPython arena and module leaks and exit non-zero.

Suggested fix: Restore ASAN_OPTIONS=detect_leaks=0 (or an LSan suppressions file for libpython) for the Python launcher path, while keeping leak detection on for native gtests. If the removal is intentional, add evidence that TheRock ASAN pytest runs stay green.

mangupta · 2026-10-15T03:31:56Z

+    start = index
+    parts = [_strip_comments(lines[index]).strip()]


🤖 AI-assisted pre-review (for @mangupta)

[IMPORTANT] Track multiline block-comment state before scanning environment variables

The source scanner can interpret text inside multiline C/C++ block comments as an environment-variable read, parameter declaration, or accessor use, causing the required RCCL env-var check to fail on comment-only changes.

Suggested fix: Tokenize each file with persistent block-comment state shared by declaration, quoted-string, and accessor scanning. Add tests where getenv("NCCL_..."), NCCL_PARAM(...), and ncclParam... appear on interior lines of /* ... */ comments both with and without a leading *, and verify that no finding is produced.

mangupta · 2026-10-15T03:31:56Z

  opcode_ = inst_.op;
 }

+void Mimg::capture_nsa_words(const MachineInst *inst, const Operand *vaddr) {
+  if (!inst_.nsa)
+    return;
+  nsa_vaddr_operand_ = vaddr;
+  size_ = sizeof(OpEncoding) + sizeof(MachineInst);


🤖 AI-assisted pre-review (for @mangupta)

[IMPORTANT] Derive the full NSA encoding length before copying

Mimg::capture_nsa_words always records exactly one additional MachineInst word for every NSA instruction, regardless of how many non-sequential-address words the instruction requires.

Suggested fix: Calculate the NSA extension-word count from the encoding's address count or the supplied vaddr operand, set size_ to the complete encoded length, and copy all required words while checking against raw_words_.size(). Add a decode/execution test using an RDNA3 NSA MIMG instruction that requires at least two extension words and verify both operand decoding and advancement to the following instruction.

mangupta · 2026-10-15T03:31:56Z

+      *binding, std::make_shared<GpuVmAccessState>(std::move(translator), std::move(memory),
+                                                   binding->access_state->fault_reporter));
  ++binding->translation_epoch;


🤖 AI-assisted pre-review (for @mangupta)

[BLOCKING] Initialize the GART access state before publishing it

The first GART publication dereferences binding->access_state even though the GART binding is created with a null access state.

Suggested fix: Create an empty GpuVmAccessState when initializing the GART binding, or make publish_gart() handle a null prior state and pass an empty fault reporter when constructing the first published generation. Apply the same invariant anywhere that dereferences a binding's access state.

mangupta · 2026-10-15T03:31:56Z

+        if (!(words[0] & (1u << (24 + context.xcc_id))))
+          skip_commands(skip);


🤖 AI-assisted pre-review (for @mangupta)

[IMPORTANT] Retire predicated-off packets from native PM4 rings

A packet skipped because its predicate is false bypasses root-ring cursor retirement and publication.

Suggested fix: Treat a predicated-off packet as successfully consumed: run the root-ring retirement/publication path before continuing. Refactor retirement into a scope guard or common helper so every successfully fetched packet, including NOP and predicated-off packets, advances the root consumer cursor exactly once.

mangupta · 2026-10-15T03:31:56Z

 * @cond @tag{gpu_bm_linux} @tag{guest_windows} @tag{host} @endcond
 */
 typedef struct {
-  char driver_version[AMDSMI_MAX_STRING_LENGTH];
-  char driver_date[AMDSMI_MAX_STRING_LENGTH];
-  char driver_name[AMDSMI_MAX_STRING_LENGTH];
+  char driver_version[AMDSMI_MAX_STRING_LENGTH];         //!< Driver version, such as
+                                                         //!< 6.19.14.31400000
+  char driver_date[AMDSMI_MAX_STRING_LENGTH];            //!< Driver release date
+  char driver_name[AMDSMI_MAX_STRING_LENGTH];            //!< Driver name
+  char driver_kernel_version[AMDSMI_MAX_STRING_LENGTH];  //!< amdgpu kernel source version,
+                                                         //!< such as 6.19.14
+  char amdgpu_driver_version[AMDSMI_MAX_STRING_LENGTH];  //!< amdgpu module version, such
+                                                         //!< as 31400000


🤖 AI-assisted pre-review (for @mangupta)

[BLOCKING] Preserve the ABI of amdsmi_driver_info_t

Appending four 256-byte arrays grows the public amdsmi_driver_info_t from 768 to 1792 bytes without a versioned replacement or reserved storage, breaking existing binary callers.

Suggested fix: Keep amdsmi_driver_info_t at its existing size and layout. Expose the additional version strings through a new versioned structure/API or another ABI-compatible extension mechanism; do not have the existing amdsmi_get_gpu_driver_info() write beyond the legacy structure.

mangupta · 2026-10-15T03:31:57Z

+// Annoyingly, unrecognized attributes are a warning and the scoped
+// atomic builtins predate clang_atomic_attributes.
+#if defined(__has_extension) && __has_extension(clang_atomic_attributes)
+#define __HIP_ATOMIC_ATTR(...) [[clang::atomic(__VA_ARGS__)]]
+#else
+#define __HIP_ATOMIC_ATTR(...)
+#endif


🤖 AI-assisted pre-review (for @mangupta)

[SUGGESTION] __HIP_ATOMIC_ATTR is defined in a public header and never undefined

amd_hip_atomic.h defines __HIP_ATOMIC_ATTR(...) at file scope with no push_macro/pop_macro or #undef. The macro leaks into every user translation unit. amd_hip_unsafe_atomics.h defines a near-identical __HIP_ATOMICS_ATTR and does scope it with push_macro/pop_macro. Two spellings of the same helper, one leaked and one scoped, are easy to confuse.

Suggested fix: Reuse a single macro name, and scope it with push_macro/pop_macro (or #undef at the end of the header), as amd_hip_unsafe_atomics.h does.

mangupta · 2026-10-15T03:31:57Z

+
+// ================================================================================================
+void StatCO::ForEachGlobalVar(void (*cb)(const void*, const char*, size_t,
+                                         const void*)) {
+  std::scoped_lock lock(sclock_);
+  // Resolved here rather than by the caller: this sweep runs from HRR's
+  // capture init, which is inside hip::init(), and re-entering the runtime
+  // through hipGetSymbolAddress from there deadlocks on the init once-flag.
+  // No device is current that early, and GetGlobalVar asserts rather than
+  // erroring on a negative device id, so the address is reported as
+  // unresolved instead.
+  // No device is current this early in init, and GetGlobalVar asserts rather
+  // than erroring on a negative id, so fall back to device 0 — which is the
+  // device a recording resolves its globals on anyway.
+  int device_id = ihipGetDevice();
+  if (device_id < 0) device_id = 0;
+  const bool resolvable = static_cast<size_t>(device_id) < g_devices.size();
+  for (const auto& [host_var, var] : vars_) {
+    if (var == nullptr) continue;
+    hipDeviceptr_t dev_ptr = nullptr;
+    size_t sym_size = 0;
+    if (resolvable &&
+        GetGlobalVar(host_var, device_id, &dev_ptr, &sym_size) != hipSuccess)
+      dev_ptr = nullptr;
+    cb(host_var, var->GetName().c_str(), var->GetSize(), dev_ptr);
+  }
+}


🤖 AI-assisted pre-review (for @mangupta)

[SUGGESTION] ForEachGlobalVar has contradictory comments and eagerly materializes every device global on device 0 during init

The first comment block says that when no device is current, "the address is reported as unresolved". The second block, and the code, fall back to device 0 and call GetGlobalVar for every registered variable. Under HRR capture, hip::init() therefore loads code objects and materializes every registered global on device 0, while holding sclock_. This happens even for processes that never use device 0, and it adds startup cost proportional to the number of globals.

Suggested fix: Delete the stale comment. Consider resolving device addresses lazily: record name and size only, and resolve at replay or on first use. If eager resolution stays, document the device-0 materialization side effect.

mangupta · 2026-10-15T03:31:57Z

+
+function(hrr_verify_generated_files)
+  set(updated_files)
+
+  foreach(kind HEADER CAPTURE PLAYBACK)
+    set(generated "${GENERATED_${kind}}")
+    set(expected "${EXPECTED_${kind}}")
+
+    if(NOT EXISTS "${generated}")
+      message(FATAL_ERROR "HRR code generation did not produce ${generated}")
+    endif()
+    if(NOT EXISTS "${expected}")
+      message(FATAL_ERROR "HRR checked-in generated file is missing: ${expected}")
+    endif()
+
+    execute_process(
+      COMMAND "${CMAKE_COMMAND}" -E compare_files --ignore-eol "${generated}" "${expected}"
+      RESULT_VARIABLE compare_result)
+    if(NOT compare_result EQUAL 0)
+      execute_process(
+        COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${generated}" "${expected}"
+        RESULT_VARIABLE copy_result)
+      if(NOT copy_result EQUAL 0)
+        message(FATAL_ERROR "Failed to update checked-in HRR generated file: ${expected}")
+      endif()
+      list(APPEND updated_files "${expected}")
+    endif()
+  endforeach()
+
+  if(updated_files)
+    string(JOIN "\n  " updated_files_text ${updated_files})
+    message(FATAL_ERROR
+      "HRR generated files were stale relative to newly generated output and were updated in the working tree:\n"
+      "  ${updated_files_text}\n\n"
+      "Add the updated files to the change and rerun the build.\n"
+      "Review whether new or changed HIP APIs require explicit classification in\n"
+      "projects/hrr/tools/gen_hrr_api_args.py, and then run script again.\n")
+  endif()


🤖 AI-assisted pre-review (for @mangupta)

[SUGGESTION] HRR codegen check rewrites checked-in source files during the build

When CLR_HRR_CODEGEN_CHECK is on, hrr_verify_generated_files copies freshly generated output over files in the source tree (projects/hrr/include, projects/clr/hipamd/src/hrr, projects/hrr/playback) and then fails the build. Those same files are listed in the custom command's DEPENDS. As a result:

a build with a read-only or shared source tree errors on the copy;

the build mutates the checkout;

the next build re-runs the rule because its inputs changed.

Suggested fix: Make the check report-only: show a diff and fail without copying. Leave regeneration to an explicit target or script (for example, a hrr_codegen_update target).

mangupta · 2026-10-15T03:31:57Z

    int status;
    //send packet to av_codec
    status = avcodec_send_packet(dec_context_, av_pkt);
-    if (status < 0) {
+    if (status == AVERROR(EAGAIN) || status == AVERROR_EOF || status == AVERROR(EINVAL)) {
+        // Benign flow control in avcodec's send/receive/flush model, not a decode error:
+        // EAGAIN = decoder is full and must be drained first; EOF = already flushed;
+        // EINVAL = a flush is required before the next packet (drain-state transition).
+        // The authoritative decode-failure signal is avcodec_receive_frame() below.
+        FunctionExitLog(g_rocdec_logger);
+        return 0;
+    }
+    else if (status < 0) {


🤖 AI-assisted pre-review (for @mangupta)

[SUGGESTION] avcodec_send_packet EAGAIN/EINVAL now returns success without draining, so the packet is dropped while ROCDEC_SUCCESS is reported

On AVERROR(EAGAIN) the new code returns 0 before reaching the avcodec_receive_frame loop. The packet is never resubmitted and buffered frames are never drained. The comment itself says the decoder "must be drained first". AVERROR(EINVAL) is also returned for genuine misuse, such as a codec that is not opened. Both cases are now reported to SubmitDecode as success.

Suggested fix: On EAGAIN, drain with avcodec_receive_frame and then resend the packet. Limit the EINVAL fast path to the flush/EOF state transition, and propagate other EINVAL cases as errors.

mangupta · 2026-10-15T03:31:57Z

+## rocm-dbgapi-0.82.1 for ROCm-X
+### Added
+- Support for the following architectures:
+ - `gfx1170`
+ - `gfx1171`
+ - `gfx1172`
+ - `gfx10-7-generic`


🤖 AI-assisted pre-review (for @mangupta)

[SUGGESTION] CHANGELOG lists gfx10-7-generic but the code adds gfx11-7-generic

The 0.82.1 changelog entry announces support for gfx10-7-generic. The code and the public header add gfx11-7-generic (EF_AMDGPU_MACH_AMDGCN_GFX11_7_GENERIC, gfx11_7_generic_t).

Suggested fix: Change gfx10-7-generic to gfx11-7-generic in CHANGELOG.md.

mangupta · 2026-10-15T05:57:57Z

🤖 AI-assisted pre-review (for @mangupta)

Re-reviewed the updated head (a780e1109f90) — no new findings at or above the posting threshold.

amd-vkale changed the title ~~[DO NOT MERGE] Remove callbacks from [Do not merge] rocProfiler-SDK: integration preview of all four services~~ [DO NOT MERGE] Remove callbacks from rocProfiler-SDK: integration preview of all four services Sep 23, 2026

github-actions Bot added documentation Improvements or additions to documentation project: rocprofiler-sdk labels Sep 23, 2026

therock-pr-bot Bot added the Not ready to Review label Sep 23, 2026

amd-vkale requested a balanced review from Copilot September 23, 2026 18:04

amd-vkale marked this pull request as ready for review September 23, 2026 18:04

amd-vkale requested review from a team as code owners September 23, 2026 18:04

Copilot started reviewing on behalf of amd-vkale September 23, 2026 18:04

View session

Copilot started work on behalf of amd-vkale September 23, 2026 18:04

View session

Copilot finished work on behalf of amd-vkale September 23, 2026 18:10

Copilot AI reviewed Sep 23, 2026

View reviewed changes

Copilot started work on behalf of amd-vkale September 23, 2026 18:11

View session

Copilot finished work on behalf of amd-vkale September 23, 2026 18:13

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

Bratislav Filipovic and others added 5 commits September 30, 2026 13:36

Merge branch 'develop' into users/anujshuk-amd/ainic-pmc-target-arch-nic

8f52971

amd-vkale and others added 15 commits October 6, 2026 20:42

mangupta reviewed Oct 15, 2026

View reviewed changes

		start = index
		parts = [_strip_comments(lines[index]).strip()]

		if (!(words[0] & (1u << (24 + context.xcc_id))))
		skip_commands(skip);

Conversation

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

Uh oh!

Issue Tracking

Uh oh!

therock-pr-bot Bot commented Sep 23, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Policy Checks Passed

Uh oh!

therock-pr-bot Bot commented Sep 23, 2026

Uh oh!

amd-vkale commented Sep 23, 2026

Uh oh!

amd-vkale commented Sep 23, 2026

Uh oh!

Copilot AI commented Sep 23, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

cursor Bot Sep 25, 2026

Choose a reason for hiding this comment

Uh oh!

cursor Bot Sep 25, 2026

Choose a reason for hiding this comment

Uh oh!

cursor Bot Sep 25, 2026

Choose a reason for hiding this comment

Uh oh!

cursor Bot Sep 25, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI commented Sep 23, 2026

Uh oh!

amd-argus Bot commented Oct 7, 2026

Uh oh!

mangupta left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

mangupta commented Oct 15, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

20 participants

amd-vkale commented Sep 23, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 23, 2026 •

edited

Loading