diff --git a/projects/rccl/src/transport/net_ib_cast/common_cast.h b/projects/rccl/src/transport/net_ib_cast/common_cast.h index 69152a03b3fd..526d02d7b905 100644 --- a/projects/rccl/src/transport/net_ib_cast/common_cast.h +++ b/projects/rccl/src/transport/net_ib_cast/common_cast.h @@ -66,7 +66,7 @@ extern union ncclSocketAddress IbCastIfAddr; enum ncclIbRequestMatchingScheme { BY_INDEX = 0, BY_ID = 1, - BY_ORDER = 2, // send requests are posted in the order they are received (CTS OFFLOAD) + BY_ORDER = 2, // completions matched by posted WR identity (CTS offload, or software CTS with on-demand recvs) }; struct ncclIbMr { @@ -705,6 +705,8 @@ ncclResult_t IbCastBaseCommInit(struct ncclIbNetCommBase* baseComm, bool isSend) ncclResult_t IbCastRecvCommInit(struct ncclIbRecvComm* recvComm); ncclResult_t IbCastSendCommInit(struct ncclIbSendComm* sendComm); +bool IbCastByOrderRequested(); + struct ncclIbListenComm { int dev; struct ncclSocket sock; @@ -821,6 +823,7 @@ extern int64_t rcclParamIbQpSchedUpdateInterval(); extern int64_t rcclParamIbQpSchedSplitDataMin(); extern int64_t rcclParamIbQpSchedLogInterval(); +void IbCastReportMatchingScheme(); #define QP_SCHED_WEIGHT_ENV_VAR "RCCL_IB_QP_SCHED_WEIGHT" #define QP_SCHED_WEIGHT_ENV_VAR_ALIAS "NCCL_IB_QP_SCHED_WEIGHT" #define QP_SCHED_LOG_PATH_ENV_VAR "RCCL_IB_QP_SCHED_LOG_PATH" diff --git a/projects/rccl/src/transport/net_ib_cast/connect.cc b/projects/rccl/src/transport/net_ib_cast/connect.cc index 83e1fb921b85..d96f03811b32 100644 --- a/projects/rccl/src/transport/net_ib_cast/connect.cc +++ b/projects/rccl/src/transport/net_ib_cast/connect.cc @@ -69,9 +69,10 @@ static inline bool IbCastIsCtsOffloadEnabled(int isP2p) { static int IbCastResolveRecvMatchingScheme(bool useCtsOffload) { // Order matters here: - // BY_ORDER -> ctsoffload - // BY_ID -> failover - // BY_INDEX -> default or user requested + // BY_ORDER -> CTS offload (forced), else AINIC/user request after remaps + // BY_ID -> failover / OOO RQ (beats explicit or default BY_ORDER) + // BY_INDEX -> default (non-AINIC) or user requested + // Explicit SCHEME=BY_ORDER is not a hard force: same as AINIC default. if (useCtsOffload) { return BY_ORDER; @@ -82,12 +83,29 @@ static int IbCastResolveRecvMatchingScheme(bool useCtsOffload) { } int64_t requested = ncclParamIbCastReceiverSideMatchingScheme(); - if (requested == -2 || requested == BY_ORDER) { - return BY_INDEX; + if (requested == -2) { + return IbCastAinicRoce ? BY_ORDER : BY_INDEX; } return requested; } +extern int64_t ncclParamIbCastReceiverSideMatchingScheme(); +// Resolve live: IbCastOffloadEnabled is still mutated after the first call +// (PORT_FAILOVER clears it in IbCastInitDevices). Caching BY_ORDER from the +// pre-failover value would disable resiliency on comms that actually run BY_ID. +bool IbCastByOrderRequested() { + return IbCastResolveRecvMatchingScheme(IbCastOffloadEnabled) == 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))); +} + + ncclResult_t IbCastInitCommDevBase(int ibDevN, struct ncclIbNetCommDevBase* base, void* cq_context, int cqSize) { base->ibDevN = ibDevN; ncclIbDev* ibDev = IbCastDevs + ibDevN; diff --git a/projects/rccl/src/transport/net_ib_cast/init.cc b/projects/rccl/src/transport/net_ib_cast/init.cc index fda9f0713c48..b12e2b7d9ce4 100644 --- a/projects/rccl/src/transport/net_ib_cast/init.cc +++ b/projects/rccl/src/transport/net_ib_cast/init.cc @@ -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))) { diff --git a/projects/rccl/src/transport/net_ib_cast/p2p.cc b/projects/rccl/src/transport/net_ib_cast/p2p.cc index 635017996de5..bc78bf55304e 100644 --- a/projects/rccl/src/transport/net_ib_cast/p2p.cc +++ b/projects/rccl/src/transport/net_ib_cast/p2p.cc @@ -16,7 +16,7 @@ NCCL_PARAM(IbCastArThreshold, "IB_AR_THRESHOLD", -2); 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). 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. - // This works out that each CTS posting QP gets drained - if (comm->useCtsOffload && (slot == ctsQp->ctsQpSlot)) { - wr.send_flags |= IBV_SEND_SIGNALED; - wr.wr_id = (req - req->base->reqs); - IbCastAddEvent(req, ctsQp->devIndex); + // 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 { - req = &(commBase->reqs[wc->wr_id]); + NCCLCHECK(IbCastRequestRetrieveAsIndex(commBase->reqs, (uint32_t)wc->wr_id, &req)); } if (req == NULL) { diff --git a/projects/rccl/src/transport/net_ib_cast/p2p_resiliency.cc b/projects/rccl/src/transport/net_ib_cast/p2p_resiliency.cc index a684c7cdddbc..8b79d3ad5936 100644 --- a/projects/rccl/src/transport/net_ib_cast/p2p_resiliency.cc +++ b/projects/rccl/src/transport/net_ib_cast/p2p_resiliency.cc @@ -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 (ncclParamIbCastResiliencyPortFailover() == 0) { + if (IbCastByOrderRequested() || ncclParamIbCastResiliencyPortFailover() == 0) { INFO(NCCL_NET, "NET/IB: %s: Resiliency is disabled on the %s communicator (comm=%p)", __func__, baseComm->isSend ? "send" : "recv", baseComm); *resCtx = NULL; diff --git a/projects/rccl/src/transport/net_ib_cast/scheduler.cc b/projects/rccl/src/transport/net_ib_cast/scheduler.cc index 4a7a827df397..8fd01aa8eb41 100644 --- a/projects/rccl/src/transport/net_ib_cast/scheduler.cc +++ b/projects/rccl/src/transport/net_ib_cast/scheduler.cc @@ -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 " + "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; if (!base->schedParmsInit) { base->schedParms = castGlobalQpSchedParms; base->schedParmsInit = true; diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json index b71b13ff2374..64cc4a01acde 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json @@ -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)", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json index b49d9e81c74d..383fdae6ae93 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json @@ -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)",