fix(rccl): default gfx120x single-node collectives to Simple protocol

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.

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;

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);
+        }
       });
 }