From 135fde0e25e17503f3b9b7b932f7e5e1e0f4a53a Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Thu, 8 Oct 2026 16:52:41 -0400 Subject: [PATCH 1/4] fix(rccl): default gfx120x single-node collectives to Simple protocol LL can deadlock single-node gfx1200/gfx1201 collectives: the LL FIFO polls in prims_ll.h use non-temporal loads without the acquire ordering that the Simple protocol uses on gfx120x. Automatic selection still picked LL for small Broadcast/Reduce/AllGather/ReduceScatter/AllReduce calls unless NCCL_P2P_DISABLE=1 was set; every collective in PyTorch's FSDP state-dict test (4 B to 60.6 KB per rank) selected LL on gfx1201. Select Simple for automatic single-node gfx120x collectives. NCCL_PROTO still overrides the choice, and multi-node behavior, including the NCCL_P2P_DISABLE override, is unchanged. JIRA ID: ROCM-29721 Refs: https://github.com/ROCm/rocm-systems/issues/5480 Co-authored-by: Cursor --- projects/rccl/src/rccl_wrap.cc | 4 ++- projects/rccl/test/host/wrap-test.cc | 45 +++++++++++++++++----------- 2 files changed, 30 insertions(+), 19 deletions(-) diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index 108cba404082..6b52c15c5141 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -181,7 +181,9 @@ 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, diff --git a/projects/rccl/test/host/wrap-test.cc b/projects/rccl/test/host/wrap-test.cc index 5bd1f8c092a8..86211926d5fb 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,48 @@ 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); + constexpr ncclFunc_t funcs[] = {ncclFuncAllGather, ncclFuncReduceScatter, ncclFuncAllReduce}; + constexpr size_t nBytes[] = {60600 * 4, 60600 * 4, 16000}; + for (const char* arch : {"gfx1200", "gfx1201"}) { + ncclComm* comm = MakeCommWithArch(arch); + comm->nNodes = 1; + comm->nRanks = 4; + for (size_t index = 0; index < sizeof(funcs) / sizeof(funcs[0]); ++index) { + SCOPED_TRACE(::testing::Message() << "arch=" << arch << " func=" << funcs[index]); + ncclTaskColl info{}; + info.func = funcs[index]; + info.protocol = NCCL_PROTO_LL; + rcclUpdateCollectiveProtocol(comm, nBytes[index], &info); + EXPECT_EQ(NCCL_PROTO_SIMPLE, info.protocol); + } + DeleteCommWithArch(comm); + } }); } From d293f5a4dc633774b5024f035f584a53429a21a3 Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Thu, 8 Oct 2026 17:02:24 -0400 Subject: [PATCH 2/4] docs(rccl): note gfx120x single-node Simple default in CHANGELOG Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 10f5cd492240..d38e62084b06 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -45,6 +45,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed tuner plugins receiving uninitialized cost-model constants. * Fixed GIN proxy descriptor shared-memory sizing and alignment. * Fixed virtual address space exhaustion when symmetric windows backed by the same physical allocation are registered repeatedly. +* 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. From c69f82b177c912a2dfb20e3c6220562d7fdb942e Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Mon, 12 Oct 2026 01:42:50 -0400 Subject: [PATCH 3/4] test(rccl): pin NCCL_PROTO on gfx120x and update the P2P-disable comment Add UpdateCollectiveProtocol_Gfx120xUserOverrideKeepsLL so NCCL_PROTO=LL must survive on single-node gfx1200/gfx1201; the gfx950 override test could not catch the gfx120x arm ignoring NCCL_PROTO. Use a struct of cases in Gfx120xSingleNodeDefaultsSimple instead of parallel arrays, and stop calling gfx120x LL functional in the NCCL_P2P_DISABLE comment, which now only matters for multi-node. JIRA ID: ROCM-29721 Co-authored-by: Cursor --- projects/rccl/src/rccl_wrap.cc | 8 ++---- projects/rccl/test/host/wrap-test.cc | 41 ++++++++++++++++++++++++---- 2 files changed, 37 insertions(+), 12 deletions(-) diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index 21c201e220f3..9c47c739b633 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -185,12 +185,8 @@ void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, s // 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 86211926d5fb..1b5f152bdc17 100644 --- a/projects/rccl/test/host/wrap-test.cc +++ b/projects/rccl/test/host/wrap-test.cc @@ -1193,18 +1193,25 @@ TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSi "Wrap_UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSimple", []() { SetMicroEnvAbsent("NCCL_PROTO"); - constexpr ncclFunc_t funcs[] = {ncclFuncAllGather, ncclFuncReduceScatter, ncclFuncAllReduce}; - constexpr size_t nBytes[] = {60600 * 4, 60600 * 4, 16000}; + // Totals, not per rank: rcclGetSizePerRank divides by nRanks for AllGather/ReduceScatter, not 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 + }; for (const char* arch : {"gfx1200", "gfx1201"}) { ncclComm* comm = MakeCommWithArch(arch); comm->nNodes = 1; comm->nRanks = 4; - for (size_t index = 0; index < sizeof(funcs) / sizeof(funcs[0]); ++index) { - SCOPED_TRACE(::testing::Message() << "arch=" << arch << " func=" << funcs[index]); + for (const auto& c : kCases) { + SCOPED_TRACE(::testing::Message() << "arch=" << arch << " func=" << c.func); ncclTaskColl info{}; - info.func = funcs[index]; + info.func = c.func; info.protocol = NCCL_PROTO_LL; - rcclUpdateCollectiveProtocol(comm, nBytes[index], &info); + rcclUpdateCollectiveProtocol(comm, c.nBytes, &info); EXPECT_EQ(NCCL_PROTO_SIMPLE, info.protocol); } DeleteCommWithArch(comm); @@ -1212,6 +1219,28 @@ TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSi }); } +// 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); + } + }); +} + TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx110xSingleNodeDelegatesToProtocolSelector) { RUN_ISOLATED_TEST( "Wrap_UpdateCollectiveProtocol_Gfx110xSingleNodeDelegatesToProtocolSelector", From 01ebf55e3c09593c2af360500d2b15fce768cbba Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Tue, 13 Oct 2026 03:32:10 -0400 Subject: [PATCH 4/4] test(rccl): pin small gfx120x collectives and note the dormant LL table Add Broadcast and Reduce cases at 4 B per rank to Gfx120xSingleNodeDefaultsSimple, so routing only the small-cutoff funcs back to rcclGetProtoForGfx120x() now fails the test. Note on rcclGetProtoForGfx120x() that it has no production caller while single-node gfx120x defaults to Simple, and is kept to re-enable LL once the prims_ll.h loads are fixed. JIRA ID: ROCM-32656 Co-authored-by: Cursor --- projects/rccl/src/rccl_wrap.cc | 2 ++ projects/rccl/test/host/wrap-test.cc | 4 +++- 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index 9c47c739b633..5c952980441c 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, diff --git a/projects/rccl/test/host/wrap-test.cc b/projects/rccl/test/host/wrap-test.cc index 1b5f152bdc17..7d96b07f6154 100644 --- a/projects/rccl/test/host/wrap-test.cc +++ b/projects/rccl/test/host/wrap-test.cc @@ -1193,7 +1193,7 @@ TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSi "Wrap_UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSimple", []() { SetMicroEnvAbsent("NCCL_PROTO"); - // Totals, not per rank: rcclGetSizePerRank divides by nRanks for AllGather/ReduceScatter, not AllReduce. + // Totals, not per rank: rcclGetSizePerRank divides by nRanks for every func here except AllReduce. constexpr struct { ncclFunc_t func; size_t nBytes; @@ -1201,6 +1201,8 @@ TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx120xSingleNodeDefaultsSi {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);