diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index bf2ae721ea6..97b4acb203b 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -84,6 +84,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed a crash (`HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`) on gfx110x and gfx120x when the Tree algorithm runs over shared memory with more channels than the Tree graph search found, for example with `NCCL_ALGO=Tree` and `NCCL_P2P_DISABLE=1`. Channels added after `ncclTopoPreset` had no intra-node tree links; `ncclTopoPostset` now sets the tree and CollNet chain links on every channel, as it already did for rings. * Fixed the host one-sided RMA API being unavailable on multi-node communicators whose ranks had fused NICs. NIC fusion merges devices for the regular transport only; GIN and RMA select from their own topology node types and never receive a fused device. * Fixed Anvil SDMA GIN small-message AllToAll latency after `ncclGinFenceLevel` support. A sub-threshold `gin.put`, a `gin.putValue`, or a standalone `gin.signal` emitted two system-scope fences per peer: one in `fenceBeforeSignal` and one in `signalPeer`. `signalPeer` now emits a single system-scope release immediately before the remote atomic, rather than an agent-scope release before it and a system-scope fence after it, which also stops payload stores from passing the signal cell. A clean-queue IPC `SignalInc` skips `fenceBeforeSignal` entirely, since it would only repeat that release. The SDMA `quiet()` before a signal now runs only when the call itself issued SDMA; previously any call carrying a signal or counter quieted the peer queue, including an IPC put after an earlier SDMA put to the same peer, which now relies on `Flush` to drain that queue. The `NCCL_GIN_ANVIL_SDMA_IPC_AGENT_FENCE` debug knob is removed. Each `Flush` thread drains the peers it owns and then clears exactly the bits it drained, over a mask disjoint from every other thread's. Previously `Flush` loaded a snapshot, quieted it, and cleared the whole word on the coop's leading thread, on every rank, which could drop a concurrent mark whose doorbell had not yet fired and then skip the quiet that mark needed. Clearing after the drain also keeps the invariant that a clear bit means a drained queue, so a second coop flushing the same context cannot return from `Flush` while the first is still inside `quiet()`. +* Fixed P2P protocol selection on single-node all-XGMI gfx942 and gfx950. The NCCL 2.29 channel estimate counted SIMPLE channels, which is 1 below 64 KiB per peer, so the 4 KiB `NCCL_P2P_LL_THRESHOLD` never selected legacy LL and transfers in that range lost the low-latency path. With the variable unset or empty, legacy LL is now used up to the smaller of the SIMPLE part size (`NCCL_P2P_NVL_CHUNKSIZE / 8`, 64 KiB by default) and 64 KiB per peer. A non-empty `NCCL_P2P_LL_THRESHOLD` still uses the per-channel rule, and the LL128 path is unchanged. Other architectures, PCIe-only pairs, and multi-node communicators are unchanged. ### Known issues diff --git a/projects/rccl/docs/userguide/source/env.rst b/projects/rccl/docs/userguide/source/env.rst index 9f579057e0b..9211cd0cd45 100644 --- a/projects/rccl/docs/userguide/source/env.rst +++ b/projects/rccl/docs/userguide/source/env.rst @@ -1945,9 +1945,11 @@ NCCL_P2P_LL_THRESHOLD The ``NCCL_P2P_LL_THRESHOLD`` is the maximum per-channel message size (in bytes) at or below which RCCL uses the legacy LL low-latency protocol for P2P (send/recv) operations. Above this size, RCCL uses the SIMPLE protocol. This threshold applies when the P2P latency path uses the legacy LL protocol (i.e. when the LL128 path is not active for the communicator). Because legacy LL carries roughly 2x wire overhead, it stops beating SIMPLE at relatively small sizes, so this defaults lower than ``NCCL_P2P_LL128_THRESHOLD``. +On single-node all-XGMI gfx942 and gfx950, an unset or empty value does not use that per-channel product. The SIMPLE channel estimate is 1 for messages up to the SIMPLE part size, so ``nChannels * 4096`` would not select LL. RCCL instead allows legacy LL up to the smaller of that part size (``NCCL_P2P_NVL_CHUNKSIZE / 8``, 64 KiB with the default 512 KiB chunk) and 64 KiB per peer. A non-empty value keeps the per-channel rule on every architecture. The LL128 path uses ``NCCL_P2P_LL128_THRESHOLD`` instead and is not widened. + Values accepted ^^^^^^^^^^^^^^^ -Decimal number. Default is 4096. +Decimal number. Default is 4096. An empty value is treated as unset. NCCL_P2P_LL128_THRESHOLD ------------------------ diff --git a/projects/rccl/src/enqueue/enqueue.cc b/projects/rccl/src/enqueue/enqueue.cc index 679a42e266e..57741267ad4 100644 --- a/projects/rccl/src/enqueue/enqueue.cc +++ b/projects/rccl/src/enqueue/enqueue.cc @@ -1291,6 +1291,12 @@ static ncclResult_t scheduleCollTasksToPlan(struct ncclComm* comm, struct ncclKe // relatively early -- internode alltoall sweeps show it regressing past ~16-32 KiB/channel. Default // 4 KiB: the largest cutoff that avoids the legacy-LL mid-size regressions while keeping its win. NCCL_PARAM(P2pLLThreshold, "P2P_LL_THRESHOLD", 4096); +// Non-empty only. ncclLoadParam treats "" as the default, so an empty value is not a user set. +// Read on each call: this is not the parameter's cached load, and a later set or unset must show. +static bool p2pLLThresholdUserSet() { + const char* str = ncclGetEnv("NCCL_P2P_LL_THRESHOLD"); + return str && str[0] != '\0'; +} // Separate, independent threshold for the LL128 latency path (used on gfx942/gfx950 with // NCCL_ALLOC_P2P_NET_LL_BUFFERS=1 and comm LL128 enabled, and on gfx1250 when // NCCL_P2P_LL128_ENABLE=1). Default 16 KiB: internode alltoall/scatter/gather sweeps (1-16 nodes) @@ -1522,7 +1528,16 @@ static ncclResult_t addP2pToPlan(struct ncclComm* comm, struct ncclKernelPlan* p } else { bool lat = useLL128OptIn ? hasLL128[dir] : hasLL[dir]; ssize_t latencyThreshold = useLL128OptIn ? ncclParamP2pLL128Threshold() : ncclParamP2pLLThreshold(); - if (!(useLL128OptIn && latencyThreshold == 0)) lat &= bytes[dir] <= nChannels[dir] * latencyThreshold; + ssize_t latencyBytesMax = nChannels[dir] * latencyThreshold; + // SIMPLE's channel count is 1 here, so the per-channel LL threshold never selects LL on + // single-node gfx942/gfx950. Allow LL up to p2pChunkSize/8, capped at 64 KiB/peer. A + // non-empty NCCL_P2P_LL_THRESHOLD, and the LL128 path, keep the stock rule. + bool const userSetP2pLLThreshold = p2pLLThresholdUserSet(); + if (comm->nNodes <= 1 && comm->isAllNvlink && !useLL128OptIn && + (comm->cudaArch == 940 || comm->cudaArch == 950) && !userSetP2pLLThreshold) { + latencyBytesMax = std::max(latencyBytesMax, std::min(comm->p2pChunkSize / 8, 65536)); + } + if (!(useLL128OptIn && latencyThreshold == 0)) lat &= bytes[dir] <= latencyBytesMax; protocol[dir] = lat ? (useLL128OptIn ? NCCL_PROTO_LL128 : NCCL_PROTO_LL) : NCCL_PROTO_SIMPLE; } } diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index 5109cc0179f..4520ed36497 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -4361,6 +4361,80 @@ struct Gfx1250AddP2pScene { return nullptr; } }; + +// Single-node all-XGMI gfx942/gfx950 legacy-LL gate. cudaArch is 100*major+10*minor, +// so gfx942 is 940; BatchPlanComm's default 942 does not enter the branch. +// p2pChunkSize is the single-node default (512 KiB), so the SIMPLE estimate is 1 +// channel at or below 64 KiB/peer. nChannelsMax is 1 on purpose: the bug is the +// protocol choice, not how many channels the LL recompute would then stripe. +struct Gfx94P2pLlOpts { + int cudaArch = 940; + int nNodes = 1; + int isAllNvlink = 1; + bool ll128OptIn = false; + int p2pChunkSize = 1 << 19; +}; + +struct Gfx94P2pLlScene { + BatchPlanComm bp; + std::unique_ptr topo{new ncclTopoSystem{}}; + std::vector peers; + std::vector peerSlots; + char llBuf[8]{}; + char ll128Buf[8]{}; + char sendMem[8]{}; + char recvMem[8]{}; + int planTotalTasks[2]{1, 1}; + + explicit Gfx94P2pLlScene(Gfx94P2pLlOpts opts) : bp(opts.nNodes, opts.cudaArch) { + InstallSendRecvDevFuncIds(); + ncclComm* c = bp.c(); + c->rank = 0; + c->nRanks = 2; + c->nNodes = opts.nNodes; + c->maxLocalRanks = 2; // nNodes > 1 divides the P2P channel base by this + c->cudaArch = opts.cudaArch; + c->isAllNvlink = opts.isAllNvlink; + c->allocP2pNetLLBuffers = opts.ll128OptIn; + c->ll128LineElems = 8; + c->ll128DataElems = 7; + c->p2pnChannels = 1; + c->p2pnChannelsPerPeer = 1; + c->p2pChannelShiftSize = 0; + c->p2pChunkSize = opts.p2pChunkSize; + c->buffSizes[NCCL_PROTO_LL] = 1 << 20; + c->buffSizes[NCCL_PROTO_LL128] = 1 << 20; + c->buffSizes[NCCL_PROTO_SIMPLE] = 1 << 20; + SetSingleGpuArch(topo.get(), opts.cudaArch == 950 ? "gfx950" : "gfx942"); + topo->ll128Enabled = opts.ll128OptIn; + c->topo = topo.get(); + + peers.resize(2); + peerSlots.resize(2); + for (int i = 0; i < 2; ++i) { + peerSlots[i] = &peers[i]; + for (int ci = 0; ci < NCCL_MAX_CONNS; ++ci) { + peers[i].send[ci].conn.buffs[NCCL_PROTO_LL] = peers[i].recv[ci].conn.buffs[NCCL_PROTO_LL] = llBuf; + peers[i].send[ci].conn.buffs[NCCL_PROTO_LL128] = peers[i].recv[ci].conn.buffs[NCCL_PROTO_LL128] = ll128Buf; + } + } + c->channels[0].peers = peerSlots.data(); + } + + ncclResult_t add(ssize_t bytes) { + struct ncclTaskP2p* tasks[2] = {nullptr, nullptr}; + return addP2pToPlan(bp.c(), bp.p(), /*nChannelsMin=*/1, /*nChannelsMax=*/1, /*p2pRound=*/0, + /*sendRank=*/1, sendMem, bytes, /*recvRank=*/1, recvMem, bytes, + /*sendOpCount=*/0, /*recvOpCount=*/0, planTotalTasks, tasks); + } + + ncclDevWorkP2p* p2pWork() { + for (auto* node = bp.p()->workQueue.head; node != nullptr; node = node->next) { + if (node->workType == ncclDevWorkTypeP2p) return reinterpret_cast(node + 1); + } + return nullptr; + } +}; } // namespace TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx1250Enable1Mixed2KiB8KiB_OneLl128WorkItem) { @@ -4553,6 +4627,98 @@ TEST_F(EnqueueMicrotest, AddPersistentPlanRef_DropsNonLiveRefsAndKeepsLiveOnes) ncclPersistentPlanRefDrop(registered); } +// Pin the numeric cutoff through the redirected ncclParamP2pLLThreshold (g_loadParam, +// no cache). Whether the user set NCCL_P2P_LL_THRESHOLD is a separate ncclGetEnv +// read: the widen applies when that name is absent or empty, so an explicit 4096 +// stays on the stock per-channel rule even though the number matches the default. +void PinP2pLlThreshold(bool userSet) { + if (userSet) SetMicroEnv("NCCL_P2P_LL_THRESHOLD", "4096"); + else SetMicroEnvAbsent("NCCL_P2P_LL_THRESHOLD"); + SetParam("P2P_LL_THRESHOLD", 4096); +} + +void ExpectGfx94RecvProtoLL(Gfx94P2pLlOpts opts, ssize_t bytes, int protoLL) { + Gfx94P2pLlScene sc(opts); + ASSERT_EQ(ncclSuccess, sc.add(bytes)) << "arch=" << opts.cudaArch << " bytes=" << bytes; + auto* work = sc.p2pWork(); + ASSERT_NE(nullptr, work) << "arch=" << opts.cudaArch << " bytes=" << bytes; + EXPECT_EQ(protoLL, int(work->recvProtoLL)) << "arch=" << opts.cudaArch << " bytes=" << bytes; + EXPECT_EQ(protoLL, int(work->sendProtoLL)) << "arch=" << opts.cudaArch << " bytes=" << bytes; +} + +void ExpectGfx94RecvProtoLL(int cudaArch, ssize_t bytes, int protoLL) { + Gfx94P2pLlOpts opts; + opts.cudaArch = cudaArch; + ExpectGfx94RecvProtoLL(opts, bytes, protoLL); +} + +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94UnsetLlThreshold_WidensLlThrough64KiB) { + PinP2pLlThreshold(/*userSet=*/false); + const struct { + ssize_t bytes; + int protoLL; + } cases[] = { + {16 << 10, 1}, {32 << 10, 1}, {64 << 10, 1}, {(64 << 10) + 1, 0}, + {128 << 10, 0}, {256 << 10, 0}, {512 << 10, 0}, + }; + for (const auto& c : cases) ExpectGfx94RecvProtoLL(940, c.bytes, c.protoLL); + // gfx950 shares the gate. 942 is not 100*9+10*4, so the default BatchPlanComm + // arch must stay SIMPLE at 16 KiB. + ExpectGfx94RecvProtoLL(950, 16 << 10, 1); + ExpectGfx94RecvProtoLL(942, 16 << 10, 0); +} + +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94UserLlThreshold_KeepsPerChannelRule) { + PinP2pLlThreshold(/*userSet=*/true); + ExpectGfx94RecvProtoLL(940, 16 << 10, 0); + ExpectGfx94RecvProtoLL(950, 16 << 10, 0); +} + +// NCCL_P2P_LL_THRESHOLD= is present but empty. ncclLoadParam uses the default, so +// the widen still applies. A nullptr-only check would treat this as user-set. +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94EmptyLlThreshold_Widens) { + SetMicroEnv("NCCL_P2P_LL_THRESHOLD", ""); + SetParam("P2P_LL_THRESHOLD", 4096); + ExpectGfx94RecvProtoLL(940, 16 << 10, 1); +} + +// 32 KiB is inside the widen (stock cutoff is 4 KiB). Each gate that is off must +// stay SIMPLE; dropping nNodes, isAllNvlink, or !useLL128OptIn would pick LL. +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94WidenGates_StaySimple) { + PinP2pLlThreshold(/*userSet=*/false); + // nNodes > 1 takes the channel base through rcclParamPxnOptQpUsage, which the + // host fake aborts on. PXN disabled is the production default and skips that + // call. The widen still requires a single node. + { + ScopedHook pxnDisable(g_pxnDisable, [](struct ncclComm*) { return 1; }); + Gfx94P2pLlOpts multiNode; + multiNode.nNodes = 2; + ExpectGfx94RecvProtoLL(multiNode, 32 << 10, 0); + } + Gfx94P2pLlOpts pcie; + pcie.isAllNvlink = 0; + ExpectGfx94RecvProtoLL(pcie, 32 << 10, 0); + Gfx94P2pLlOpts ll128; + ll128.ll128OptIn = true; + ExpectGfx94RecvProtoLL(ll128, 32 << 10, 0); +} + +// The default 512 KiB chunk makes p2pChunkSize/8 equal the 64 KiB cap, so either +// operand of the min would pass. A smaller chunk stops at chunk/8. A larger one +// stays at the cap. +TEST_F(EnqueueMicrotest, AddP2pToPlan_Gfx94ChunkSize_BoundsTheWiden) { + PinP2pLlThreshold(/*userSet=*/false); + Gfx94P2pLlOpts small; + small.p2pChunkSize = 256 << 10; + ExpectGfx94RecvProtoLL(small, 32 << 10, 1); + ExpectGfx94RecvProtoLL(small, (32 << 10) + 1, 0); + Gfx94P2pLlOpts large; + large.p2pChunkSize = 1 << 20; + ExpectGfx94RecvProtoLL(large, 64 << 10, 1); + ExpectGfx94RecvProtoLL(large, (64 << 10) + 1, 0); + ExpectGfx94RecvProtoLL(large, 128 << 10, 0); +} + // =========================================================================== // rawTaskAppend / ncclCeInitJob (the NCCL_ENQUEUE_REARCH_ENABLE=1 path), then argument validation arms. // ===========================================================================