Repository navigation
fix(rccl): fixed udma balancing for p2p comms (develop) - #12948
Conversation
Undocumented RCCL default-behavior changes
Paste into
Describe these changes in
|
✅ All Policy Checks Passed
📖 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 |
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: IB_UDMA_POLICY defaults to 0, so the p2p balancing this PR is named for only arms when a user opts in, and nothing in the tree sets it.
Comment @argus review to re-run, or @argus re-review after pushing.
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: connect.cc does not compile at this head. Line 83 redefines ncclIbCastUBMALow instead of defining ncclIbCastUBMAHigh, and line 85 returns the deleted type ncclIbCastUBMA.
Comment @argus review to re-run, or @argus re-review after pushing.
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: ncclIbCreateQpIonic now holds the general device mutex across ibv_create_qp, so AINIC QP creation serializes against MR registration on the default path.
Comment @argus review to re-run, or @argus re-review after pushing.
alex-breslow-amd
left a comment
alex-breslow-amd
left a comment
There was a problem hiding this comment.
You have a lot of duplicate code with attribute creation spread across multiple files and functions. This should get cleaned up.
|
@karthikarum Since you were commenting on my PR re: this issue.. #12275 what you think about this PR, kindly take a look as well? thanks! |
Hi @sunguobao |
The previous commit left connect.cc uncompilable: the high-pipeline constant was declared under the low name and ncclIbCastSelectUDMA still returned the removed enum type. Fix both, rename the UBMA identifiers to UDMA to match the rest of the feature, and warn at init when RCCL_IB_UDMA_POLICY is outside 0-3 instead of silently falling back to policy 0. Co-authored-by: Cursor <cursoragent@cursor.com>
- Name the RCCL_IB_UDMA_POLICY values with enum struct ncclIbCastUDMAPolicy (First/Last bound the init range check) and document that QP sharing overrides channel affinity. - Drop qpIndexInDev in favour of qpIdx. A new IbCastQpStripe helper computes devIndex and qpIdx and fails with ncclInternalError on a zero device count; AINIC port recovery restores qpIdx after IbCastQpCreateAttrInitSharing resets it. - Guard mask selection and ibv_create_qp with a dedicated per-device uDMA mutex instead of ncclIbDev::mutex, so QP creation no longer serializes against MR registration. - Warn once per device when ionic_dv_pd_set_udma_mask fails. - Log the active uDMA policy at init. Co-Authored-By: Claude <noreply@anthropic.com>
accaf7d to
5909ad9
Compare
|
@argus re-review |
amd-arozanov
left a comment
amd-arozanov
left a comment
There was a problem hiding this comment.
The PR description is outdated after ba423a0: qpIndexInDev was removed in favor of qpIdx and IbCastQpStripe, and QP creation now uses the dedicated nccl_udma_mutex instead of the general device mutex. Please update these two bullets before merge?
alex-breslow-amd
left a comment
alex-breslow-amd
left a comment
There was a problem hiding this comment.
Two questions about whether the opt-in policies balance the MI355X single-QP alltoall config, plus two nits. The Test Result section still says performance numbers will come after the runs. Since the default stays at 0 and the opt-in policies are the whole fix, could you add those numbers before merge?
TheRock Submodule Bump ActivityNewest first
|
## Motivation Follow-up to #12948 (`RCCL_IB_UDMA_POLICY`), addressing the review comments that were left after it merged: documentation for the new variable, a guard on the policy enum range, one shared validity check, a consistency check in AINIC QP recovery, and the QP-count requirement of `SplitConnection`. ## Technical Details One commit per review item: - **Count sentinel for `ncclIbCastUDMAPolicy`** (`common_cast.h`): `Last` is now `Count - 1`, with `static_assert`s that `First == ChannelAffinity == 0` and `Last == SplitConnection`. Adding a policy without updating the asserts fails to compile. Nothing `switch`es over the enum, so `Count` cannot trigger `-Wswitch`. - **Shared range check**: `ncclIbCastUDMAPolicyIsValid()` in `common_cast.h`, used by both `init.cc` (WARN + fallback) and `ncclIbCastGetUDMAPolicy()` in `connect.cc`. No behavior change. - **Recovery `devIndex` check** (`p2p_resiliency_recovery_ainic.cc`): the device index from `IbCastQpStripe` must match `localQp->devIndex`; otherwise WARN and return `ncclInternalError`. Connect sets `devIndex` from the same stripe, so this cannot fire today; it catches the striping assumption behind `qpIdxInDev` breaking later. - **`SplitConnection` QP-count note** (comments only): the default `NCCL_IB_QPS_PER_CONNECTION` of 2 already splits a connection. With one QP per device per connection (`NCCL_IB_QPS_PER_CONNECTION=1`, or `RCCL_IB_QPS_PER_P2P=1` for P2P) `qpIdx` is always 0 and the policy reduces to channel affinity by channel-id popcount parity. RMA comms also use one QP. - **Docs**: `RCCL_IB_UDMA_POLICY` entry in `docs/api-reference/env-variables.rst` and a CHANGELOG "Added" bullet. Policy 0 pins each channel's send QPs, and separately its receive QPs, to one pipeline per device. Policy 1 round-robins P2P QPs on one per-device counter. Policy 3's default of two QPs per connection is enough; splitting is lost with one QP per connection and on RMA comms. Default behavior (`RCCL_IB_UDMA_POLICY` unset = 0) is unchanged. A per-direction round-robin counter for policy 1 was tried and dropped: it did not improve performance. ## Issue Tracking JIRA ID: ROCM-31118 ## Test Plan - Compile check of the changed translation units (`connect.cc`, `init.cc`, `p2p_resiliency_recovery_ainic.cc`, and through them `common_cast.h`) on develop `df181d33333`. - Negative compile check: adding a policy value before `Count` without updating the asserts fails the build. - `rccl_check_default_behavior.py --base-ref <develop>`: passes. ## Test Result - Compile: 0 errors, no new warnings in the touched files. - Env-var check: "No undocumented RCCL or NCCL environment-variable changes." For reference, #12948's validation on TW (MI355X + AINIC, 8 reps, interleaved arms, 144/144 runs clean per stage) ran with the default 2 QPs per connection, where this PR does not change the assignment. Policy unset was flat vs develop on alltoall/all_reduce/all_gather at 2N and 4N. alltoall Avg busbw: 2N +9.3% (policy 1, 8/8), 4N +6.0% (policy 1, 8/8); policies 2 and 3 +3.6% at 4N. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
Motivation
Multi-node alltoall on MI355X with AI-NIC regresses at all message sizes in the August RCCL code drop. The root cause is uneven use of the two AINIC uDMA pipelines by P2P QPs.
[internal-env]
Technical Details
ncclIbCreateQpIonicintoncclIbCastSelectUDMAand addsRCCL_IB_UDMA_POLICY:0(default): current behavior. With QP sharing, QPs alternate by sharing group (or by QP index when there is one group); otherwise every QP of a channel is pinned to one pipeline, assigned alternately as channels are first seen.1: P2P QPs round-robin across pipelines in creation order; collectives keep the policy 0 behavior.2: the PD mask is left unchanged, so the driver picks the pipeline.3: each connection is split across both pipelines; the channel id picks the starting pipeline and consecutive per-device QPs alternate.qpIdxandIbCastQpStripetoncclIbQpCreateAttr, set on the sender, receiver and AINIC port-recovery QP creation paths.nccl_udma_mutexfrom pipeline selection throughibv_create_qp, because the PD and its uDMA mask are shared by all comms on the device.RCCL_IB_UDMA_POLICYis outside 0-3; such values behave as 0.The default stays 0 because the new policies regress other workloads, so they are opt-in until they have been evaluated with the customer. This is the develop counterpart of #12390, which targets
rccl/drop_2026-08.Issue Tracking
JIRA ID: ROCM-31118
Test Plan
alltoall_perfon MI355X + AI-NIC with eachRCCL_IB_UDMA_POLICYvalue.Test Result
The compile check passes.
BASE1, BASE2 (both develop 5909ad9), PR (ba423a0)
2 nodes results (mi355+AINIC):
Alltoall (out-of-place busbw, GB/s; deltas are against base1 matched repeat by repeat, with wins out of 8)
All_reduce and all_gather (regression check runs)
4 nodes results (mi355+AINIC):
Alltoall (out-of-place busbw, GB/s; deltas are against base1 matched repeat by repeat, with wins out of 8)
Allreduce/Allgather is flat as for 2N.
Submission Checklist
(AI Asist)