fix(rccl): net-ib-cast: enable by_order for no cts-offload path (drop 2026 08) by alexander-sannikov · Pull Request #11172 · ROCm/rocm-systems

alexander-sannikov · 2026-09-03T10:28:11Z

Motivation

This PR enables BY_ORDER on the no-CTS-offload path and makes it the default on AINIC when the matching scheme is unset, so users do not need to set NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=2 manually for the common AINIC + RCCL_CTS_OFFLOAD_ENABLED=0 configuration.

Technical Details

Changes are confined to net_ib_cast:

IbCastResolveRecvMatchingScheme() (connect.cc) — centralizes scheme selection:
- CTS offload enabled → BY_ORDER (unchanged)
- Explicit NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=2 → BY_ORDER
- Unset default (-2) on AINIC → BY_ORDER (new)
- Explicit =0 → BY_INDEX (user override preserved)
IbCastByOrderRequested() — now reflects the resolved scheme (CTS offload, AINIC default, or explicit env), not only an explicit =2.
Incompatible features disabled when BY_ORDER is active: QP scheduler, split-data, WRR, resiliency/port-failover, and out-of-order RQ (via scheduler.cc, init.cc, p2p_resiliency.cc).
IbCastReportMatchingScheme() — init-time INFO logging of collective vs P2P CTS offload state and resolved RecvMatchingScheme.

Issue Tracking

JIRA ID: AICOMNET-396

Test Plan

Defined in Jira

Test Result

Provided in Jira

Submission Checklist

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

… env values Run the whole of net_ib_transport.json and mi300x_mellanox_ib.json over the path PR ROCm#11172 enables, instead of only the four dedicated suites: - Global env now carries NCCL_NET=IB-CAST and NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=2, so every suite that does not set its own value runs BY_ORDER over the CAST plugin. - CTS offload is selected by --system: the "ainic" profile pins RCCL_CTS_OFFLOAD_ENABLED=0, which is the case this change is about, and the new "ainic_cts" profile is the same environment with it set to 1. Suites keep their own value where they set one. - cast_base pins the matching scheme back to auto. Its suites assert the behaviour of the QP scheduler, WRR and split-data, all of which BY_ORDER turns off, so under the global setting they would assert against a scheduler that is not running. Verified: those suites report recvMatchingScheme=0 and pass, while suites inheriting the global setting report recvMatchingScheme=2 with useCtsOffload=0. Quote environment values when building non-MPI test commands. The AINIC profile sets NCCL_NET_FORCE_MERGE=rdma0,rdma1;rdma2,rdma3, and unquoted the ';' ends the shell command, so the rest was executed as a program: "/bin/sh: 1: rdma2,rdma3: not found". That failed every single-rank test on AINIC, Socket_BootstrapPollNoBusyLoop among them.

nikitaxgusev

one comment, LGTM

Copilot

🟡 Changes recommended

Configuration caching and incomplete feature gating can disable resiliency or leave incompatible receive and scheduling paths active.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

Enables BY_ORDER receive matching for AINIC without CTS offload and updates related transport behavior.

Changes:

Centralizes receive-matching scheme resolution and reporting.
Adds software BY_ORDER completion handling.
Attempts to disable incompatible scheduling, resiliency, and OOO-RQ features.

File summaries

File	Description
`connect.cc`	Resolves and reports matching schemes.
`p2p.cc`	Handles software BY_ORDER completions and CTS signaling.
`scheduler.cc`	Gates scheduling for BY_ORDER.
`p2p_resiliency.cc`	Disables resiliency for BY_ORDER.
`init.cc`	Warns about incompatible features and reports configuration.
`common_cast.h`	Updates declarations and matching documentation.

Review details

Files reviewed: 6/6 changed files
Comments generated: 6
Review effort level: Balanced

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

sunguobao · 2026-09-04T22:47:09Z

+      if (ncclParamIbCastOooRq()) {
+	      WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d):"
+              " disabling out-of-order RQ", BY_ORDER);


Same question here - about the mechanism of disabling OOO.

Separately, I see prepost was touched in earlier commits in this PR but later reverted - does prepost need a force-disable too, in this case?

Also, I see prepostReceiveWorkRequests changes in older commits but now gone. Is this intended (to leave prepost behavior untouched even if BY_ORDER is picked)?

+static int IBCastUsedMatchingScheme = -1;
+bool IbCastByOrderRequested() {
+  if (IBCastUsedMatchingScheme < 0) {
+    IBCastUsedMatchingScheme = IbCastResolveRecvMatchingScheme(IbCastOffloadEnabled);
+  }
+  return IBCastUsedMatchingScheme == BY_ORDER;
+}


sunguobao · 2026-09-05T00:17:05Z

    return enabled;
  }

+  if (IbCastByOrderRequested()) {


sorta agree with copilot here - any consideration between putting this at it current place, or before line 27? I am not having all the implications known, so just a genuine question.

 }

 void IbCastUpdateSchedParmsTry(struct ncclIbNetCommBase* base, int nreqs, int size) {
+  if (base->recvMatchingScheme == BY_ORDER) return;


+              " disabling out-of-order RQ", BY_ORDER);
+      }
+    }
+    IbCastReportMatchingScheme();


 int64_t IbCastArThreshold = 8192;

-// By default, use ncclIbRequestMatchingScheme::BY_INDEX matching scheme.
+// Default -2: BY_INDEX. 0=BY_INDEX, 1=BY_ID, 2=BY_ORDER (software CTS; CTS offload forces BY_ORDER).


amd-argus

ARGUS

Findings

1 finding(s) that fall outside the diff and could not be anchored to a line

projects/rccl/CHANGELOG.md:13 (Low) No CHANGELOG entry for the user-visible AINIC default change

Bottom line: Worth settling before approval: BY_ORDER receive only works because the split-data disable at scheduler.cc:190 is ineffective, and the cast tests already drive the sender to one QP through ncclIbCastSetSchedParms.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus · 2026-09-11T06:50:49Z

 }

 void IbCastUpdateSchedParmsTry(struct ncclIbNetCommBase* base, int nreqs, int size) {
+  if (base->recvMatchingScheme == BY_ORDER) return;


This early return is what currently keeps BY_ORDER working, which is the opposite of what the warning at scheduler.cc:36 claims it does. The receiver arms one WQE per QP across base->nqps (p2p.cc:630 and 633-649) and matches by receive-queue FIFO position, and IbCastCompletionEventByOrder has no equivalent of the single-QP compensation that the BY_INDEX handler does at p2p.cc:1095-1106, so any smaller sender fan-out leaves IbCastRequestIsComplete (p2p.cc:814-820) unsatisfied. ncclIbCastSetSchedParms (scheduler.cc:492) already reaches that state: CastTests.cpp:737, :855 and :887 set schedEnable true with splitData false, which hits oneQp at scheduler.cc:381 and returns nqps 1. Could the invariant be made explicit before the split-data disable is made effective?

amd-argus · 2026-09-11T06:50:49Z

-  if (requested == -2 || requested == BY_ORDER) {
-    return BY_INDEX;
+  if (requested == -2) {
+    return IbCastAinicRoce ? BY_ORDER : BY_INDEX;


Defaulting to BY_ORDER here makes IbCastUpdateSchedParmsTry return at scheduler.cc:190 before it copies castGlobalQpSchedParms, so base->schedParms stays zeroed on every cast comm. ncclIbCastGetSchedState (scheduler.cc:453-456) then reports schedEnable false, doWrr false and splitDataMin 0, while CastSchedParmsReflectEnvVars (CastTests.cpp:331-333) asserts the opposite under cast_base (mi455_ainic_roce_func.json:499-511), which never sets the env var. The token-consumption tests in CastTests.cpp go the same way. Could the cast configs pin NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=0, or the tests move with the default?

amd-argus · 2026-09-11T06:50:50Z

  }

+  if (IbCastByOrderRequested()) {
+    WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d): disabling the QP "


This warning fires on the plain AINIC default path where the user set nothing, and it fires twice per init: once through init.cc:612 and once through IbCastQpSchedInitParms at scheduler.cc:78, which init.cc:642 always calls. On AINIC the scheduler is already disabled by default at scheduler.cc:48-52, so nothing new is being turned off here. Could this be an INFO, emitted once from the init path rather than from inside the predicate?

amd-argus · 2026-09-11T06:50:50Z

+extern int64_t ncclParamIbCastReceiverSideMatchingScheme();
+static int IBCastUsedMatchingScheme = -1;
+bool IbCastByOrderRequested() {
+  if (IBCastUsedMatchingScheme < 0) {


The IBCastUsedMatchingScheme < 0 sentinel collides with values the resolver can return, since connect.cc:92 passes requested through verbatim and only -2 is special-cased at connect.cc:89. With NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=-1 the value never latches and the predicate re-resolves on every call. A separate static bool for the latch would be sturdier. While here, IBCast is the only occurrence of that prefix in the tree, everything else is IbCast.

amd-argus · 2026-09-11T06:50:50Z

 extern int64_t rcclParamIbQpSchedSplitDataMin();
 extern int64_t rcclParamIbQpSchedLogInterval();

+void IbCastReportMatchingScheme();


IbCastReportMatchingScheme leaves the resolved scheme observable only through INFO, so nothing can assert which scheme a comm actually picked. ncclIbCastSchedState (net_ib_cast_inspect.h:27-39) is already filled by ncclIbCastGetSchedState (scheduler.cc:435-459) from the same ncclIbNetCommBase that carries recvMatchingScheme (common_cast.h:505), so one int field there would make the default testable. That header is outside this PR, so this is a follow-up rather than a change to make here.

amd-argus · 2026-09-11T06:50:50Z

+      }
+      if (ncclParamIbCastOooRq()) {
+	      WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d):"
+              " disabling out-of-order RQ", BY_ORDER);


Suggested change

" disabling out-of-order RQ", BY_ORDER);

WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d):"

" disabling out-of-order RQ", BY_ORDER);

amd-argus · 2026-09-11T06:50:50Z

  return requested;
 }

+extern int64_t ncclParamIbCastReceiverSideMatchingScheme();


Suggested change

extern int64_t ncclParamIbCastReceiverSideMatchingScheme();

amd-argus · 2026-09-11T06:50:50Z

+       IbCastIsCtsOffloadEnabled(0)? "ON":"OFF" , IbCastResolveRecvMatchingScheme(IbCastIsCtsOffloadEnabled(0)));
+  INFO(NCCL_NET, "NET/IB-CAST: P2P communicators: CTS Offload: %s   RecvMatchingScheme: %d",
+       IbCastIsCtsOffloadEnabled(1)? "ON":"OFF" , IbCastResolveRecvMatchingScheme(IbCastIsCtsOffloadEnabled(1)));
+}


Suggested change

}

void IbCastReportMatchingScheme() {

INFO(NCCL_NET, "NET/IB-CAST: collectives communicators: CTS Offload: %s RecvMatchingScheme: %d",

IbCastIsCtsOffloadEnabled(0) ? "ON" : "OFF", IbCastResolveRecvMatchingScheme(IbCastIsCtsOffloadEnabled(0)));

INFO(NCCL_NET, "NET/IB-CAST: P2P communicators: CTS Offload: %s RecvMatchingScheme: %d",

IbCastIsCtsOffloadEnabled(1) ? "ON" : "OFF", IbCastResolveRecvMatchingScheme(IbCastIsCtsOffloadEnabled(1)));

}

github-actions Bot added the project: rccl label Sep 3, 2026

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

alexander-sannikov requested review from VadimKutovoi, ddurnov and nikitaxgusev September 3, 2026 15:31

nikitaxgusev approved these changes Sep 3, 2026

View reviewed changes

Comment thread projects/rccl/src/transport/net_ib_cast/init.cc Resolved

VadimKutovoi approved these changes Sep 3, 2026

View reviewed changes

alexander-sannikov added the ainic-test label Sep 3, 2026

alexander-sannikov marked this pull request as ready for review September 3, 2026 16:15

alexander-sannikov requested a review from a team as a code owner September 3, 2026 16:15

i-kosarev reviewed Sep 3, 2026

View reviewed changes

Comment thread projects/rccl/src/transport/net_ib_cast/p2p_resiliency.cc Resolved

i-kosarev requested a balanced review from Copilot September 3, 2026 18:34

Copilot started reviewing on behalf of i-kosarev September 3, 2026 18:35

View session

Copilot AI reviewed Sep 3, 2026

View reviewed changes

amd-argus Bot reviewed Sep 11, 2026

View reviewed changes

alexander-sannikov changed the title ~~fix(rccl): net-ib-cast: enable by_order for no cts-offload path~~ fix(rccl): net-ib-cast: enable by_order for no cts-offload path (drop 2026 08) Sep 21, 2026

alexander-sannikov added 6 commits September 21, 2026 13:40

fix(rccl): net-ib-cast: enable by_order for no cts-offload path

7d6e624

By-order: minimal code cleanup

5a2c302

Automatic selection of by-order on AINIC

e8510c1

Improved debug output

a8c91bc

Added BY_order into test configs

b5c8b49

BY-ORDER: fixed review comments

c6177c0

alexander-sannikov force-pushed the users/asanniko/by_order branch from d12b0dd to c6177c0 Compare September 21, 2026 12:44

alexander-sannikov merged commit 4401caa into rccl/drop_2026-08 Sep 21, 2026
93 of 98 checks passed

alexander-sannikov deleted the users/asanniko/by_order branch September 21, 2026 18:44

	" disabling out-of-order RQ", BY_ORDER);
	WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d):"
	" disabling out-of-order RQ", BY_ORDER);

-}
+void IbCastReportMatchingScheme() {
+  INFO(NCCL_NET, "NET/IB-CAST: collectives communicators: CTS Offload: %s   RecvMatchingScheme: %d",
+       IbCastIsCtsOffloadEnabled(0) ? "ON" : "OFF", IbCastResolveRecvMatchingScheme(IbCastIsCtsOffloadEnabled(0)));
+  INFO(NCCL_NET, "NET/IB-CAST: P2P communicators: CTS Offload: %s   RecvMatchingScheme: %d",
+       IbCastIsCtsOffloadEnabled(1) ? "ON" : "OFF", IbCastResolveRecvMatchingScheme(IbCastIsCtsOffloadEnabled(1)));
+}

Conversation

alexander-sannikov commented Sep 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Issue Tracking

Test Plan

Test Result

Submission Checklist

Uh oh!

nikitaxgusev left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

sunguobao Sep 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

sunguobao Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

sunguobao Sep 5, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

alexander-sannikov commented Sep 3, 2026 •

edited

Loading

sunguobao Sep 4, 2026 •

edited

Loading