Repository navigation
fix(rccl): net-ib-cast: enable by_order for no cts-offload path (drop 2026 08) #11172
There are no files selected for viewing
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
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -644,6 +644,21 @@ ncclResult_t IbCastInitDevices(ncclDebugLogger_t logFunction, ncclProfilerCallba | |
| IbCastUseInline = false; | ||
| IbCastAinicCtsInlineData = false; | ||
| } | ||
| // After all offload/scheduler mutations so BY_ORDER + failover logs match | ||
| // the scheme comms will actually use (IbCastByOrderRequested reads live). | ||
| if (ret == ncclSuccess && IbCastNDevs > 0) { | ||
| if (IbCastByOrderRequested()) { | ||
| if (ncclParamIbCastResiliencyPortFailover() > 0) { | ||
| WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d): disabling resiliency " | ||
| "(port failover and port recovery)", BY_ORDER); | ||
| } | ||
| if (ncclParamIbCastOooRq()) { | ||
| WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d):" | ||
| " disabling out-of-order RQ", BY_ORDER); | ||
| } | ||
| } | ||
| IbCastReportMatchingScheme(); | ||
| } | ||
| return ret; | ||
| fail: | ||
| if (devices && (ncclSuccess != wrap_ibv_free_device_list(devices))) { | ||
|
|
||
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -16,7 +16,7 @@ | |
| NCCL_PARAM(IbCastArThreshold, "IB_AR_THRESHOLD", -2); | ||
| int64_t IbCastArThreshold = 8192; | ||
|
|
||
| // Default -2: BY_INDEX. 0=BY_INDEX, 1=BY_ID, 2=BY_ORDER (software CTS; CTS offload forces BY_ORDER). | ||
|
|
||
| NCCL_PARAM(IbCastReceiverSideMatchingScheme, "IB_RECEIVER_SIDE_MATCHING_SCHEME", -2); | ||
| RCCL_PARAM(IbCastGdrFlushGpuMemNoRelaxedOrdering, "GDR_FLUSH_GPU_MEM_NO_RELAXED_ORDERING", 1); | ||
|
|
||
|
|
@@ -549,11 +549,15 @@ ncclResult_t IbCastPostFifo(struct ncclIbRecvComm* comm, struct ncclIbRequest* r | |
| // - The status of all posted Send Request is considered unknown | ||
| // | ||
| // slot == devIndex - When writing to CTS FIFO slot N, and this QP lives on device index N, it should send signalled. | ||
| // BY_ORDER signals only the drain slot and does not carry `|| resiliency`: | ||
| // IbCastResiliencyInit nulls the resiliency context for BY_ORDER comms. | ||
| if (comm->base.recvMatchingScheme == BY_ORDER) { | ||
| bool signalCts = comm->useCtsOffload ? (slot == ctsQp->ctsQpSlot) : (slot == ctsQp->devIndex); | ||
| if (signalCts) { | ||
| wr.send_flags |= IBV_SEND_SIGNALED; | ||
| wr.wr_id = (req - req->base->reqs); | ||
| IbCastAddEvent(req, ctsQp->devIndex); | ||
| } | ||
| } else if (!comm->useCtsOffload && (slot == ctsQp->devIndex || comm->base.resiliency)) { | ||
| wr.send_flags |= IBV_SEND_SIGNALED; | ||
| wr.wr_id = slot; | ||
|
|
@@ -898,8 +902,11 @@ static ncclResult_t IbCastCompletionEventByOrder(struct ncclIbNetCommBase* commB | |
| if (commBase->isSend) { | ||
| struct ncclIbSendComm* sendComm = (struct ncclIbSendComm*)commBase; | ||
| req = sendComm->sendReqs[wc->wr_id & 0xff][0]; | ||
| } else if (wc->opcode == IBV_WC_RDMA_READ) { | ||
| NCCLCHECK(IbCastRequestRetrieveAsIndex(commBase->reqs, (uint32_t)(wc->wr_id - NCCL_IB_FLUSH_REQ_WR_ID_OFFSET), | ||
| &req)); | ||
| } else { | ||
| NCCLCHECK(IbCastRequestRetrieveAsIndex(commBase->reqs, (uint32_t)wc->wr_id, &req)); | ||
| } | ||
|
|
||
| if (req == NULL) { | ||
|
|
||
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -598,7 +598,7 @@ static ncclResult_t IbCastResiliencyProbeProgress(struct ncclIbResiliencySend* s | |
| ncclResult_t IbCastResiliencyInit(struct ncclIbNetCommBase* baseComm, struct ncclIbResiliency** resCtx) { | ||
| assert(baseComm != NULL); | ||
| assert(resCtx != NULL); | ||
| if (IbCastByOrderRequested() || ncclParamIbCastResiliencyPortFailover() == 0) { | ||
|
alexander-sannikov marked this conversation as resolved.
Resolved
|
||
| INFO(NCCL_NET, "NET/IB: %s: Resiliency is disabled on the %s communicator (comm=%p)", __func__, | ||
| baseComm->isSend ? "send" : "recv", baseComm); | ||
| *resCtx = NULL; | ||
|
|
||
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -23,15 +23,32 @@ FILE* IbCastQpSchedLogStream; | |
| // 1. RCCL_IB_QP_SCHED_ENABLE / NCCL_IB_QP_SCHED_ENABLE explicitly set (!=−1) → honour it. | ||
| // 2. NCCL_NET=ib-cast explicitly set → force enable. | ||
| // 3. Otherwise → disable on AINIC (default path), enable everywhere else. | ||
| // BY_ORDER comms no-op WRR/split-data (IbCastQpSchedDoWrr returns immediately). | ||
| // An explicit ENABLE=1 still wins, but that combination is unsupported: warn so it | ||
| // is visible instead of silently running a scheduler that never remaps. | ||
| bool rcclUseIbCastQpSched() { | ||
| int64_t schedParam = rcclParamIbCastQpSchedEnable(); | ||
| const bool byOrder = IbCastByOrderRequested(); | ||
| if (schedParam != -1) { | ||
| bool enabled = (schedParam != 0); | ||
| INFO(NCCL_NET | NCCL_ENV, "(IB-CAST) NCCL_IB_QP_SCHED_ENABLE explicitly set to %s", | ||
| enabled ? "enabled" : "disabled"); | ||
| if (enabled && byOrder) { | ||
| WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d) but the QP " | ||
| "scheduler is force-enabled (NCCL_IB_QP_SCHED_ENABLE=%ld): unsupported combination; " | ||
| "WRR/split-data are no-ops on BY_ORDER comms", | ||
| BY_ORDER, schedParam); | ||
| } | ||
| return enabled; | ||
| } | ||
|
|
||
| if (byOrder) { | ||
| WARN("NET/IB: BY_ORDER matching requested (NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=%d): disabling the QP " | ||
|
Contributor
There was a problem hiding this comment. 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 |
||
| "scheduler, split-data and WRR", | ||
| BY_ORDER); | ||
| return false; | ||
| } | ||
|
|
||
| const char* netEnv = ncclGetEnv("NCCL_NET"); | ||
| if (netEnv && strcasecmp(netEnv, "ib-cast") == 0) { | ||
| INFO(NCCL_NET | NCCL_ENV, "(IB-CAST) NCCL_NET=ib-cast: forcing QP scheduler enabled"); | ||
|
|
@@ -180,6 +197,7 @@ void IbCastLogSched(struct ncclIbSendComm* comm) { | |
| } | ||
|
|
||
| void IbCastUpdateSchedParmsTry(struct ncclIbNetCommBase* base, int nreqs, int size) { | ||
| if (base->recvMatchingScheme == BY_ORDER) return; | ||
|
Contributor
There was a problem hiding this comment. 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 |
||
| if (!base->schedParmsInit) { | ||
| base->schedParms = castGlobalQpSchedParms; | ||
| base->schedParmsInit = true; | ||
|
|
||
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -1117,6 +1117,32 @@ | |
| "NCCL_NET": "ib-cast" | ||
| } | ||
| }, | ||
| "net_ib_transfer_cast_by_order_no_cts": { | ||
| "extends": "net_ib_transfer_cast", | ||
| "description": "ib-cast data path with software BY_ORDER matching and CTS offload forced off (the no-offload BY_ORDER arm). QP scheduler left off so TX fan-out stays at base nqps.", | ||
| "env_variables": { | ||
| "RCCL_CTS_OFFLOAD_ENABLED": "0", | ||
| "NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME": "2" | ||
| } | ||
| }, | ||
| "net_ib_multirecv_cast_by_order_no_cts": { | ||
| "extends": "net_ib_multirecv_cast", | ||
| "description": "ib-cast irecv batching under software BY_ORDER with CTS offload off — completion matching by posted WR identity rather than CTS slot.", | ||
| "env_variables": { | ||
| "RCCL_CTS_OFFLOAD_ENABLED": "0", | ||
| "NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME": "2", | ||
| "RCCL_IB_QP_SCHED_ENABLE": "0" | ||
| } | ||
| }, | ||
| "net_ib_stress_multi_qp_nosplit_cast_by_order_no_cts": { | ||
| "extends": "net_ib_stress_multi_qp_nosplit_cast", | ||
| "description": "ib-cast multi-QP (QPS=4, splitData=0) send/recv under software BY_ORDER with CTS offload off. Exercises FIFO-order matching across QPs without WRR reducing the sender QP count.", | ||
| "env_variables": { | ||
| "RCCL_CTS_OFFLOAD_ENABLED": "0", | ||
| "NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME": "2", | ||
| "RCCL_IB_QP_SCHED_ENABLE": "0" | ||
| } | ||
| }, | ||
| "unit_tests_fixtures": { | ||
| "is_gtest": true, | ||
| "binary": "rccl-UnitTestsFixtures", | ||
|
|
@@ -3785,6 +3811,12 @@ | |
| "config": "net_ib_transfer_cast", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Data Transfer (ib-cast, BY_ORDER, CTS offload off)", | ||
| "description": "ib-cast SimpleSendRecv / multiple sizes / large / outstanding transfers with NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=2 and RCCL_CTS_OFFLOAD_ENABLED=0", | ||
| "config": "net_ib_transfer_cast_by_order_no_cts", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Coverage Env (ib-cast)", | ||
| "description": "Coverage-env transfer tests via ib-cast plugin: exercises CAST GID auto-select and relaxed-ordering paths (NCCL_NET=ib-cast)", | ||
|
|
@@ -3797,6 +3829,12 @@ | |
| "config": "net_ib_multirecv_cast", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - MultiRecv (ib-cast, BY_ORDER, CTS offload off)", | ||
| "description": "ib-cast MultiRecv under software BY_ORDER with CTS offload disabled", | ||
| "config": "net_ib_multirecv_cast_by_order_no_cts", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - CTS Depth Stress (ib-cast)", | ||
| "description": "InfiniBand CTS depth stress via ib-cast plugin (NCCL_NET=ib-cast)", | ||
|
|
@@ -3833,6 +3871,12 @@ | |
| "config": "net_ib_stress_multi_qp_nosplit_cast", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Stress Multi-QP NoSplit (ib-cast, BY_ORDER, CTS offload off)", | ||
| "description": "ib-cast multi-QP no-split stress with NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=2 and RCCL_CTS_OFFLOAD_ENABLED=0", | ||
| "config": "net_ib_stress_multi_qp_nosplit_cast_by_order_no_cts", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Stress Multi-QP Fan-In (ib-cast)", | ||
| "description": "InfiniBand multi-QP fan-in stress via ib-cast plugin (NCCL_NET=ib-cast)", | ||
|
|
||
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -1102,6 +1102,32 @@ | |
| "NCCL_NET": "ib-cast" | ||
| } | ||
| }, | ||
| "net_ib_transfer_cast_by_order_no_cts": { | ||
| "extends": "net_ib_transfer_cast", | ||
| "description": "ib-cast data path with software BY_ORDER matching and CTS offload forced off (the no-offload BY_ORDER arm). QP scheduler left off so TX fan-out stays at base nqps.", | ||
| "env_variables": { | ||
| "RCCL_CTS_OFFLOAD_ENABLED": "0", | ||
| "NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME": "2" | ||
| } | ||
| }, | ||
| "net_ib_multirecv_cast_by_order_no_cts": { | ||
| "extends": "net_ib_multirecv_cast", | ||
| "description": "ib-cast irecv batching under software BY_ORDER with CTS offload off — completion matching by posted WR identity rather than CTS slot.", | ||
| "env_variables": { | ||
| "RCCL_CTS_OFFLOAD_ENABLED": "0", | ||
| "NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME": "2", | ||
| "RCCL_IB_QP_SCHED_ENABLE": "0" | ||
| } | ||
| }, | ||
| "net_ib_stress_multi_qp_nosplit_cast_by_order_no_cts": { | ||
| "extends": "net_ib_stress_multi_qp_nosplit_cast", | ||
| "description": "ib-cast multi-QP (QPS=4, splitData=0) send/recv under software BY_ORDER with CTS offload off. Exercises FIFO-order matching across QPs without WRR reducing the sender QP count.", | ||
| "env_variables": { | ||
| "RCCL_CTS_OFFLOAD_ENABLED": "0", | ||
| "NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME": "2", | ||
| "RCCL_IB_QP_SCHED_ENABLE": "0" | ||
| } | ||
| }, | ||
| "unit_tests_fixtures": { | ||
| "is_gtest": true, | ||
| "binary": "rccl-UnitTestsFixtures", | ||
|
|
@@ -2950,6 +2976,12 @@ | |
| "config": "net_ib_transfer_cast", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Data Transfer (ib-cast, BY_ORDER, CTS offload off)", | ||
| "description": "ib-cast SimpleSendRecv / multiple sizes / large / outstanding transfers with NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=2 and RCCL_CTS_OFFLOAD_ENABLED=0", | ||
| "config": "net_ib_transfer_cast_by_order_no_cts", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Coverage Env (ib-cast)", | ||
| "description": "Coverage-env transfer tests via ib-cast plugin: exercises CAST GID auto-select and relaxed-ordering paths (NCCL_NET=ib-cast)", | ||
|
|
@@ -2962,6 +2994,12 @@ | |
| "config": "net_ib_multirecv_cast", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - MultiRecv (ib-cast, BY_ORDER, CTS offload off)", | ||
| "description": "ib-cast MultiRecv under software BY_ORDER with CTS offload disabled", | ||
| "config": "net_ib_multirecv_cast_by_order_no_cts", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - CTS Depth Stress (ib-cast)", | ||
| "description": "InfiniBand CTS depth stress via ib-cast plugin (NCCL_NET=ib-cast)", | ||
|
|
@@ -2998,6 +3036,12 @@ | |
| "config": "net_ib_stress_multi_qp_nosplit_cast", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Stress Multi-QP NoSplit (ib-cast, BY_ORDER, CTS offload off)", | ||
| "description": "ib-cast multi-QP no-split stress with NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=2 and RCCL_CTS_OFFLOAD_ENABLED=0", | ||
| "config": "net_ib_stress_multi_qp_nosplit_cast_by_order_no_cts", | ||
| "enabled": true | ||
| }, | ||
| { | ||
| "name": "NET IB - Stress Multi-QP Fan-In (ib-cast)", | ||
| "description": "InfiniBand multi-QP fan-in stress via ib-cast plugin (NCCL_NET=ib-cast)", | ||
|
|
||
Oops, something went wrong.
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.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
IbCastReportMatchingSchemeleaves 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 byncclIbCastGetSchedState(scheduler.cc:435-459) from the samencclIbNetCommBasethat carriesrecvMatchingScheme(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.