Repository navigation
feat(rccl): IB cast plugin - replace UDMA idx hardcoding with runtime discovery via ionic_dv_ctx_get_udma_count - #12749
Merged
Conversation
karthikarum
requested review from
a team,
VadimKutovoi and
alexander-sannikov
and
a balanced review from Copilot
October 3, 2026 20:18
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Unconditional Ionic capability initialization corrupts the overlapping MLX5 Data Direct union member, and the tests do not exercise the production round-robin path.
Review effort: Balanced
Findings: 1
Open (4)
What changed in this PR
Adds runtime AINIC UDMA discovery and dynamic QP pinning, with host and hardware validation.
Changes:
- Loads optional Ionic UDMA discovery/readback symbols.
- Replaces two-engine pinning with runtime round-robin assignment.
- Adds unit, MPI hardware, and test-runner coverage.
| File | Description |
|---|---|
tools/scripts/test_runner/configs/mi300_ainic_ccs_roce.json |
Registers UDMA test suites. |
test/transport/NetIbMPI/CastTests.cpp |
Adds hardware pinning test. |
test/transport/NetIbCast/NetIbCastUdmaTests.cpp |
Adds host unit tests. |
test/test_categories_fixtures_debug.yaml |
Categorizes new tests. |
test/CMakeLists.txt |
Builds UDMA tests. |
src/transport/net_ib_cast/net_ib_cast_inspect.h |
Exposes test introspection APIs. |
src/transport/net_ib_cast/init.cc |
Discovers UDMA count. |
src/transport/net_ib_cast/connect.cc |
Implements dynamic pinning and readback. |
src/transport/net_ib_cast/common_cast.h |
Stores Ionic capabilities. |
src/misc/ionicdvwrap.cc |
Adds safe UDMA wrappers. |
src/misc/ionicdvsymbols.cc |
Loads optional Ionic symbols. |
src/include/ionic/ionicdvwrap.h |
Declares wrapper APIs. |
src/include/ionic/ionicdvsymbols.h |
Adds symbol function pointers. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
karthikarum
changed the title
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: on an AINIC run the new ionic.udmaCount write makes capsProvider.mlx5.dataDirect read non-zero for every device, flipping forceFlush and the DMA-BUF registration branch.
Comment @argus review to re-run, or @argus re-review after pushing.
karthikarum
force-pushed
the
users/karthik/synctoionicverbs
branch
from
October 4, 2026 02:18
0311a2d to
e59c819
Compare
Contributor
There was a problem hiding this comment.
ARGUS
Findings
2 finding(s) that could not be posted as inline comments
projects/rccl/tools/scripts/test_runner/configs/mi300_ainic_ccs_roce.json:778(Medium) cast_udma_pinning_sharing never reaches the new qpIdx sub-branch because IB_COMM_NGROUPS defaults to 4 The newudmaSelector = qpIdxsub-branch at connect.cc:491 runs only whenrcclParamIbCastCommNGroups() == 1, and qp_sharing.cc:14 declaresRCCL_PARAM(IbCastCommNGroups, "IB_COMM_NGROUPS", 4). Nothing in this config chain sets it, so the suite always takes the pre-existingqpSharingGroupIdxpath and the per-QP alternation this PR adds gets no coverage while A24 reports green. Could the suite pin the group count as well?projects/rccl/src/transport/net_ib_cast/connect.cc:484(Low) udmaCount reaches two divisors unnormalised, which becomes a SIGFPE once the init.cc write is provider-gatedudmaCountis read raw here and then used as a divisor twice, at line 493 and insidencclIbCastUdmaRoundRobinMask, neither of which guards zero. It is safe today only because init.cc:514 writes a normalised value into every device slot; once that write is restricted to non-mlx5 devices, mlx5 slots hold 0 andncclIbCreateQpIonicstill runs for them, sinceuseIonicis the same global flag rather than a per-device property. Normalising at the read site makes both divisors safe whichever way the init side lands.
Bottom line: Worth settling before approval: the capsProvider union write at init.cc:514 is still not gated on provider, so the new IB_PROVIDER_MLX5 guards do not stop an mlx5 device in AINIC mode reading back as data-direct.
Comment @argus review to re-run, or @argus re-review after pushing.
karthikarum
changed the title
karthikarum
force-pushed
the
users/karthik/synctoionicverbs
branch
from
October 13, 2026 01:25
7b6c35f to
3818dfe
Compare
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: the new EXPECT_GT(distinctEngines.size(), 2u) in CastUdmaPinningSharingVerify can never hold, because the sharing selector is one value per comm, so the sharing path still has no working check.
Comment @argus review to re-run, or @argus re-review after pushing.
VadimKutovoi
approved these changes
Oct 13, 2026
VadimKutovoi
left a comment
VadimKutovoi
left a comment
Contributor
There was a problem hiding this comment.
LGTM. Should we add the tests to other configuration files?
karthikarum
force-pushed
the
users/karthik/synctoionicverbs
branch
from
October 14, 2026 16:21
6fe87c5 to
414a80f
Compare
Contributor
Author
The test is added to ainic mi300 cfg and validated on mi300 cluster. About adding to other cluster config, I'm not sure what is available in CI/CD now and how it gets validated. If you are looking to add to a specific config file, let me know and I can follow it up in next PR. |
wenkaidu
approved these changes
Oct 15, 2026
karthikarum
force-pushed
the
users/karthik/synctoionicverbs
branch
from
October 19, 2026 02:50
414a80f to
38b7822
Compare
karthikarum
force-pushed
the
users/karthik/synctoionicverbs
branch
from
October 21, 2026 02:05
38b7822 to
28c9b1f
Compare
…iscovery Add soft-fail symbol loading for ionic_dv_ctx_get_udma_count and ionic_dv_qp_get_udma_idx in the DV wrapper layer. Store the discovered UDMA engine count in capsProvider.ionic.udmaCount (normalized: 0 falls back to 2, clamped to 8 max). Log udmaCount at device init. Guard capsProvider union reads with provider type checks (IB_PROVIDER_MLX5 for mlx5.dataDirect) in init.cc and reg.cc to prevent union corruption when the ionic union member is populated. Add shared helpers (ncclIbCastNormalizeUdmaCount, ncclIbCastUdmaRoundRobinMask) and UDMA introspection structs (ncclIbCastUdmaState, ncclIbCastTestUdmaRoundRobin, ncclIbCastTestGetDevUdmaCount) to net_ib_cast_inspect.h. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Replace hardcoded IONIC_UDMA_MASK_LOW/HIGH (2-engine assumption) in
ncclIbCastSelectUDMA() with dynamic udmaCount from capsProvider.ionic.
All policy paths (ChannelAffinity, P2pRoundRobin, SplitConnection, and
QP-sharing) now use 1u << (idx % udmaCount) for mask computation.
Fix SplitConnection policy to use (parity + qpIdx) % udmaCount instead
of parity ^ (qpIdx & 1). The old XOR-with-bit0 formula only ever
produced engine indices 0 or 1 regardless of udmaCount, starving engines
2..N-1 on hardware with more than 2 UDMA pipelines. The new additive
modulo walks all N engines while remaining algebraically identical for
the udmaCount=2 case ((a + b) mod 2 == a XOR b when a,b in {0,1}), so
there is no behavioral change on current 2-engine hardware.
Widen ncclChannelToUd.udId and nccl_channel_last_ud from bool to uint8_t
to support engine indices beyond 0/1. The channel-affinity fallback uses
ncclIbCastUdmaRoundRobinMask() (shared with unit tests via inspect.h).
Add test inspect implementations: ncclIbCastTestUdmaRoundRobin,
ncclIbCastTestGetDevUdmaCount, and ncclIbCastGetUdmaState for per-QP
UDMA index readback via ionic_dv_qp_get_udma_idx.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…inning Host-only unit tests (A22 series) for the UDMA round-robin algorithm via ncclIbCastTestUdmaRoundRobin inspect API: 1/2/4/8 engines, engine coverage, sharing-path masks, backward compat, normalize helper, and invalid args. Two-node MPI functional tests (A23): CastUdmaPinningVerify validates end-to-end UDMA pinning on AINIC hardware via ionic_dv_qp_get_udma_idx readback. CastUdmaPinningSharingVerify exercises the QP-sharing path with NGROUPS=1 (per-QP qpIdx selector). Both include data-integrity send/recv and skip cleanly on non-AINIC or older drivers. Test runner config entries added for all new tests. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
karthikarum
force-pushed
the
users/karthik/synctoionicverbs
branch
from
October 21, 2026 05:19
28c9b1f to
c09abbf
Compare
wenkaidu
approved these changes
Oct 21, 2026
Contributor
Author
|
Seeing segmentation fault with single node RCCL Coco tests which is unrelated to this change in net-ib-cast plugin. @wenkaidu @ddurnov can you pls confirm this is not related to the plugin changes.
|
VadimKutovoi
approved these changes
Oct 22, 2026
Contributor
|
@karthikarum The all-zero NET_COUNTER_TABLE. As I can see, this is a single node, 1 GPU per rank reduce_perf, so the collective runs over intra-node transport and no RoCE traffic hits the bnxt_re NICs. The collector reports deltas (before/after snapshots), so with no traffic every counter is legitimately 0. The segfault is unrelated to the plugin change. It happens at teardown: the run completes, then every rank faults in hipMemUnmap. It seems this issue: https://amd-hub.atlassian.net/browse/ROCM-33022 |
aduforest
added a commit
that referenced
this pull request
Oct 22, 2026
…unt (pre-merge hardening) #12749 made init.cc call wrap_ionicdv_ctx_get_udma_count, so the Init binary no longer linked. ionicdvwrap_fakes gains hooks for it and for wrap_ionicdv_qp_get_udma_idx, defaulting to what the real wrappers return without libionic. Tests cover the normalized per-device count and the mlx5-only data-direct checks in GetProperties and DMA-BUF registration. Co-Authored-By: Claude <noreply@anthropic.com>
aduforest
added a commit
that referenced
this pull request
Oct 22, 2026
…me count (pre-merge hardening) #12749 makes connect.cc pick uDMA engines modulo the device's discovered count, which the AINIC tests left at 0, so they died on a divide by zero. The fixture now sets the count as IbCastInitDevices would. New tests cover every policy with 3 and 4 engines, the fixed 2 without AINIC, and the uDMA inspect helpers. ionicdvwrap_fakes gains the two new uDMA hooks. Co-Authored-By: Claude <noreply@anthropic.com>
This was referenced Oct 22, 2026
pvallem
pushed a commit
that referenced
this pull request
Oct 25, 2026
…GIN (#13840) JIRA ID: AICOMRCCL-2813 ## What Host-only microtests for `src/transport/net_ib_cast/{init,common,capability,reg,gdr,multiplane,gin,net_ib_cast_gid_inspect}.cc`, in a new `rccl-UnitTestsMicroNetIbCastInit` with one test TU per unit. ## How - The verbs, mlx5/ionic, socket and neighbouring net_ib_cast seams are shared fakes. The ones the other net_ib_cast PRs also use are byte-identical copies; this PR owns the scheduler, p2p_resiliency_recovery, gdr_peermem and gdr_probe fakes. Fake ibv contexts route post_send/poll_cq back to the test. `xml.cc`, `utils.cc` and `net_telemetry.cc` are the real ones. The GDR answers latched by `std::call_once` are tested in isolated processes. - `transport_stubs.cc` gains net.cc's `NET_OPTIONAL_RECV_COMPLETION` param, and `os_fakes.cc` gains `ncclOsStrSep`, the same hunk as #13601. - Pins 10 production bugs as `KnownBug_` tests: - AINIC fuses NICs by default, although the code means to require `NCCL_IB_MERGE_NICS=1` (AICOMRCCL-3088). - A GDR probe that could not run is latched as "no GDR" (AICOMRCCL-3089). - A data-direct lookup error skips the cleanup (AICOMRCCL-3090). - A failed device init turns the next one into a silent no-op (AICOMRCCL-3091). - Deregistering unsorts the MR cache, so a cached range gets registered twice (AICOMRCCL-3092). - A regMr failure on a later device keeps the earlier devices' MRs (AICOMRCCL-3093). - GIN connect, context and registration failures leak comms and MRs (AICOMRCCL-3094). - The RMA proxy's error report indexes devices by completion slot, and can crash (AICOMRCCL-3095). - IPutSignal with data but no destination posts an empty WR instead of the signal (AICOMRCCL-3096). - A plane-index failure frees the ibv device list twice (AICOMRCCL-3143). - Not pinned, as the outcome depends on uninitialized memory: `IbCastMakeVDevice` compares and logs PCI roots it never set when a path has none (AICOMRCCL-3097). ## Coverage Lines: init.cc 642/652, common.cc 185/185, common_cast.h 57/57, capability.cc 162/162, reg.cc 116/117, gdr.cc 73/73, multiplane.cc 186/189, gin.cc 491/491 (the GDAKI arm is compiled out, as in the shipping build), net_ib_cast_gid_inspect.cc 102/102. The misses are calloc failures, a capability probe that cannot fail, a telemetry slot that is never out of range, and the return after reg.cc's loop. ## Mutation Testing 172 mutations on the hipified sources, all killed; 5 of them needed a sharper test first. 1 more against the review changes, also killed. 11 more for #12749's runtime uDMA count and its mlx5-only data-direct checks, all killed. ## Test Plan Standalone build: `rccl-UnitTestsMicroNetIbCastInit` 275/275 in about 0.4 s, and `--gtest_shuffle --gtest_repeat=20` is clean. Every other host binary was rebuilt against the changed fakes and run once. DdaColl and Rocmwrap fail as they do on develop (the break #13819 fixes) and link none of the changed files. --------- Co-authored-by: Claude <noreply@anthropic.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
AINIC hardware ships with varying numbers of UDMA engines per NIC (currently 2, but future SKUs may expose more). The net_ib_cast transport layer previously hardcoded a 2-engine assumption
(IONIC_UDMA_MASK_LOW/IONIC_UDMA_MASK_HIGH), which would silently under-utilize DMA bandwidth on hardware with more engines and could break on hardware with fewer. This PR replaces the hardcoded assumption
with runtime discovery via ionic_dv_ctx_get_udma_count and adds comprehensive test coverage — both host-only unit tests and a 2-node HW functional test — to prevent regressions and validate the pinning
end-to-end.
Technical Details
Commit 1: Dynamic UDMA count discovery (b138401)
Commit 2: Unit tests + HW functional test (b22a618)
(udmaCount=0 → fallback), and invalid argument handling
ionic_dv_qp_get_udma_idx), validates indices are in-range and consistent, then confirms data integrity with send/recv. Skips cleanly on non-AINIC or older drivers.
Issue Tracking
JIRA ID: AICOMNET-421
Test Plan
Test Result
Full test suite run on MI300 AINIC cluster (gfx942, smc300x-ccs-aus-gpucc5a + smc300x-ccs-aus-gpuf273):
A22 unit tests: all 8 cases PASSED. A23 HW test: PASSED — udmaCount=2, all QP UDMA indices in range and consistent, data integrity verified.
RCCL_4N_comparison_baseline-tot_vs_udma-get-changes_20261014.html
Submission Checklist