diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index e4a4e01b67f0..73ebfb455f2a 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -58,6 +58,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed virtual address space exhaustion when symmetric windows backed by the same physical allocation are registered repeatedly. * Fixed Copy Engine `ncclAlltoAllv` hanging on intra-node communicators of more than 8 ranks with uneven send sizes. Each rank decided from its own send sizes whether to add intra-batch barriers, so ranks could issue different barrier counts. The decision now uses the full send-size matrix, which every rank shares. * Fixed `NCCL_P2P_DISABLE` and `NCCL_P2P_LEVEL` being ignored on Intel, ARM and Zhaoxin hosts, where the CPU-based default P2P level overwrote the user setting. On such hosts `NCCL_P2P_DISABLE=1` now falls back to SHM as documented, and `NCCL_P2P_LEVEL` is honored. +* Fixed hangs in small single-node collectives on gfx1200 and gfx1201 by selecting the Simple protocol instead of LL by default. `NCCL_PROTO` still overrides this choice. ### Known issues * The FP8 ReduceSum and ReduceCopy device APIs are not exported in the LLVM bitcode library. diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index b5bedb119526..c7ea89f6b55d 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -121,6 +121,8 @@ static inline bool rcclCollSupportsRing(ncclFunc_t func) { func == ncclFuncBroadcast || func == ncclFuncReduce); } +// No production caller while single-node gfx120x defaults to Simple (see rcclUpdateCollectiveProtocol). Kept, with +// its tests, to re-enable LL once the prims_ll.h loads are fixed for gfx120x. int32_t rcclGetProtoForGfx120x(ncclFunc_t collectiveFunc, size_t sizePerRank) { int returnVal = NCCL_PROTO_SIMPLE; int SingleNodeLLCutoffs[] = {/*ncclFuncBroadcast*/ 1536, @@ -239,14 +241,12 @@ void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, s info->protocol = NCCL_PROTO_LL; } else if (!userProtocolInput && IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx120" /*match gfx120x*/)) { if (comm->nNodes == 1) { - info->protocol = rcclGetProtoForGfx120x(info->func, sizePerRank); + // LL can deadlock single-node gfx120x collectives: its FIFO polls in prims_ll.h lack the + // gfx120x acquire ordering that Simple uses. Keep Simple until LL is fixed for gfx120x. + info->protocol = NCCL_PROTO_SIMPLE; } - /** - * We prefer simple protocol when p2p_disabled = 1, - * This is due to a fix in LL protocol implementation - * for gfx120x with __MEMORY_SCOPE_SYSTEM in prims_ll.h - * causing poor performance but keeps the LL protocol functional - */ + // With P2P disabled, multi-node collectives also use Simple: the system-scope LL stores for gfx120x + // in prims_ll.h are slow. bool p2p_disabled = ncclParamP2pDisable(); if (p2p_disabled) { info->protocol = NCCL_PROTO_SIMPLE; diff --git a/projects/rccl/test/host/wrap-test.cc b/projects/rccl/test/host/wrap-test.cc index a7a9bd319451..1d0fc42bf976 100644 --- a/projects/rccl/test/host/wrap-test.cc +++ b/projects/rccl/test/host/wrap-test.cc @@ -1040,7 +1040,7 @@ TEST(WrapMicrotest, SymkHostRedOpToDev_UnknownOpReturnsNegativeOne) { // Covers the top-level user-override gate, one arch/size LL-threshold arm // (gfx950 AllGather, plus its own boundary; the gfx950/gfx942 ReduceScatter // arms right below it are the same shape with different constants and are -// covered by their own tests here), the gfx120x delegation + +// covered by their own tests here), the gfx120x single-node Simple default + // NCCL_P2P_DISABLE override, and the // nNodes>=2 minMaxLLRange-driven arm including its warn-once // undefined-tuning fallback. @@ -1167,39 +1167,79 @@ TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_UserOverrideLeavesProtocolU }); } -TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xDelegatesThenP2pDisableForcesSimple) { +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xMultiNodeP2pDisableForcesSimple) { RUN_ISOLATED_TEST( - "Wrap_UpdateCollectiveProtocol_Gfx120xDelegatesThenP2pDisableForcesSimple", + "Wrap_UpdateCollectiveProtocol_Gfx120xMultiNodeP2pDisableForcesSimple", []() { SetMicroEnvAbsent("NCCL_PROTO"); ScopedHook p2pDisable(g_paramP2pDisable, []() { return 1; }); ncclComm* comm = MakeCommWithArch("gfx1200"); - comm->nNodes = 1; - comm->nRanks = 1; + comm->nNodes = 2; + comm->nRanks = 2; ncclTaskColl info{}; info.func = ncclFuncAllGather; - info.protocol = NCCL_PROTO_LL; // whatever rcclGetProtoForGfx120x would pick, P2P_DISABLE overrides it + info.protocol = NCCL_PROTO_LL; // multi-node gfx120x gets no other tuning; only P2P_DISABLE changes it rcclUpdateCollectiveProtocol(comm, /*nBytes=*/1024, &info); EXPECT_EQ(NCCL_PROTO_SIMPLE, info.protocol); DeleteCommWithArch(comm); }); } -TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xSingleNodeDelegatesToProtocolSelector) { +// LL can deadlock single-node gfx120x collectives, so automatic selection must +// return Simple even at sizes rcclGetProtoForGfx120x maps to LL. 60,600 B per +// rank is the largest AllGather/ReduceScatter in PyTorch's FSDP state-dict test. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSimple) { RUN_ISOLATED_TEST( - "Wrap_UpdateCollectiveProtocol_Gfx120xSingleNodeDelegatesToProtocolSelector", + "Wrap_UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSimple", []() { SetMicroEnvAbsent("NCCL_PROTO"); - SetMicroEnvAbsent("NCCL_P2P_DISABLE"); - ncclComm* comm = MakeCommWithArch("gfx1200"); - comm->nNodes = 1; - comm->nRanks = 1; - ncclTaskColl info{}; - info.func = ncclFuncAllGather; - info.protocol = NCCL_PROTO_SIMPLE; - rcclUpdateCollectiveProtocol(comm, /*nBytes=*/1024, &info); - EXPECT_EQ(NCCL_PROTO_LL, info.protocol); - DeleteCommWithArch(comm); + // Totals, not per rank: rcclGetSizePerRank divides by nRanks for every func here except AllReduce. + constexpr struct { + ncclFunc_t func; + size_t nBytes; + } kCases[] = { + {ncclFuncAllGather, 60600 * 4}, // sizePerRank 60600 <= 98304 + {ncclFuncReduceScatter, 60600 * 4}, // sizePerRank 60600 <= 98304 + {ncclFuncAllReduce, 16000}, // sizePerRank 16000 <= 16384 + {ncclFuncBroadcast, 4 * 4}, // sizePerRank 4 <= 1536 + {ncclFuncReduce, 4 * 4}, // sizePerRank 4 <= 8192 + }; + for (const char* arch : {"gfx1200", "gfx1201"}) { + ncclComm* comm = MakeCommWithArch(arch); + comm->nNodes = 1; + comm->nRanks = 4; + for (const auto& c : kCases) { + SCOPED_TRACE(::testing::Message() << "arch=" << arch << " func=" << c.func); + ncclTaskColl info{}; + info.func = c.func; + info.protocol = NCCL_PROTO_LL; + rcclUpdateCollectiveProtocol(comm, c.nBytes, &info); + EXPECT_EQ(NCCL_PROTO_SIMPLE, info.protocol); + } + DeleteCommWithArch(comm); + } + }); +} + +// The gfx950 override test above can't catch the gfx120x arm ignoring +// NCCL_PROTO, so pin that the user's LL survives on single-node gfx120x too. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xUserOverrideKeepsLL) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx120xUserOverrideKeepsLL", + []() { + SetMicroEnv("NCCL_PROTO", "LL"); + for (const char* arch : {"gfx1200", "gfx1201"}) { + SCOPED_TRACE(arch); + ncclComm* comm = MakeCommWithArch(arch); + comm->nNodes = 1; + comm->nRanks = 4; + ncclTaskColl info{}; + info.func = ncclFuncAllGather; + info.protocol = NCCL_PROTO_LL; + rcclUpdateCollectiveProtocol(comm, /*nBytes=*/60600 * 4, &info); + EXPECT_EQ(NCCL_PROTO_LL, info.protocol); + DeleteCommWithArch(comm); + } }); }