fix(rccl/gin): reject Anvil GIN peers outside the LSA team

projects/rccl/CHANGELOG.md

@@ -15,6 +15,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 ### Resolved issues
 * Fixed the device linker dropping feature suffixes such as `xnack+` from `GPU_TARGETS`, which left address sanitizer builds without instrumented device code.
 * Fixed `NCCL_CUMEM_ENABLE=1` being ignored on Linux kernels older than 6.8 (for example 6.6 on MI308X). cuMem stayed off, so RCCL buffers bypassed the memory manager and `ncclCommSuspend` released no GPU memory.
+* Fixed the GIN Anvil SDMA backend (`NCCL_GIN_TYPE=7`, or `NCCL_GIN_TYPE` unset when Anvil is the selected backend) publishing signal addresses outside the mapped LSA window for GIN peers that are not in the rank's LSA team, for example an MI300X CPX communicator spanning two OAMs with `NCCL_LSA_TEAM_SIZE=4`, or any `NCCL_GIN_CONNECTION_RAIL` team with more than one rank. Such peers have no address in the window, so binding Anvil signals now fails with `ncclInvalidUsage` and a warning naming the peer, instead of warning and continuing. `ncclDevCommCreate` calls that request GIN in these configurations previously succeeded and now fail; set `NCCL_LSA_TEAM_SIZE` to cover the GIN team or select another GIN type. The symmetric GIN kernels, which address rail peers outside the LSA team, are no longer used with Anvil SDMA, so communicator initialization and collectives are unaffected.
 
 ### Known issues
 

projects/rccl/src/gin/gin_plugin_anvil_sdma.cc

@@ -785,16 +785,34 @@ static uintptr_t* ginAnvilBuildSignalPeerAddrs(void* base, int nRanks, int ginRa
   return addrs;
 }
 
+// The signal peer table is indexed by GIN peer, but each entry is a slot in this
+// rank's LSA window. Returns the first GIN peer that is not the LSA team member in
+// the slot its entry points at, or -1 if every entry is a mapped peer.
+static int ginAnvilFirstSignalPeerOutsideLsa(const struct ncclComm* comm, int nRanks, int ginRank) {
+  const struct ncclDevrState* devr = &comm->devrState;
+  const long long ginStride =
+    comm->globalGinSupport == NCCL_GIN_CONNECTION_RAIL ? comm->contiguousRanksPerHost : 1;
+  for (int pe = 0; pe < nRanks; pe++) {
+    const int slot = devr->lsaSelf + (pe - ginRank);
+    const long long worldPeer = comm->rank + (long long)(pe - ginRank) * ginStride;
+    if (devr->lsaRankList == nullptr || slot < 0 || slot >= devr->lsaSize || devr->lsaRankList[slot] != worldPeer) {
+      return pe;
+    }
+  }
+  return -1;
+}
+
 static ncclResult_t ginAnvilRegisterSignalSpan(struct ncclComm* comm, void* spanLsaSelf, size_t spanBytes, int nRanks,
                                                int ginRank) {
   struct ncclDevrState* devr = &comm->devrState;
 
-  if (nRanks != devr->lsaSize) {
-    WARN("GIN anvil-sdma: signal nRanks=%d != devr->lsaSize=%d (rank %d)", nRanks, devr->lsaSize, ginRank);
-  }
-  if (ginRank != devr->lsaSelf) {
-    WARN("GIN anvil-sdma: ginRank=%d != devr->lsaSelf=%d (signal peer indexing may be wrong)", ginRank,
-         devr->lsaSelf);
+  const int outsidePeer = ginAnvilFirstSignalPeerOutsideLsa(comm, nRanks, ginRank);
+  if (outsidePeer >= 0) {
+    WARN("GIN anvil-sdma: GIN peer %d is outside the LSA team of rank %d (ginRank=%d, ginRanks=%d, lsaSelf=%d, "
+         "lsaSize=%d). Anvil SDMA can only signal LSA team peers; set NCCL_LSA_TEAM_SIZE to cover the GIN team "
+         "or select another NCCL_GIN_TYPE.",
+         outsidePeer, comm->rank, ginRank, nRanks, devr->lsaSelf, devr->lsaSize);
+    return ncclInvalidUsage;
   }
 
   uintptr_t* hostAddrs = ginAnvilBuildSignalPeerAddrs(spanLsaSelf, nRanks, ginRank, (ptrdiff_t)devr->bigSize);

projects/rccl/src/include/sym_kernels.h

@@ -110,6 +110,7 @@ struct ncclSymkDevComm {
 struct ncclSymkState {
   bool initialized;
   bool hasLsaMultimem;
+  bool ginBackendUnsupported;
   int maxGinInboxBlocks;
   struct ncclSymkDevComm kcomm;
 };

projects/rccl/src/sym_kernels.cc

@@ -140,6 +140,13 @@ static uint32_t kernelMask_coll(ncclFunc_t coll) {
 
 NCCL_PARAM(SymGinKernelsEnable, "SYM_GIN_KERNELS_ENABLE", 1)
 NCCL_PARAM(SymRsGinChunkSize, "SYM_RS_GIN_CHUNK_SIZE", -1)
+
+// [RCCL] Anvil SDMA reaches a GIN peer only through its slot in this rank's LSA window, and every
+// rail peer of the GIN kernels is outside the LSA team.
+static bool ncclSymkGinBackendUnsupported(struct ncclComm* comm) {
+  ncclGinType_t ginType = NCCL_GIN_TYPE_NONE;
+  return ncclGetRailedGinType(comm, &ginType) == ncclSuccess && ginType == NCCL_GIN_TYPE_ANVIL_SDMA;
+}
 // [RCCL] These kernels stage tiles through a DMA engine: TMA on NVIDIA, the Tensor
 // Data Mover on gfx1250. On by default as in NCCL 2.31; ncclSymkTmaAvailable() still
 // confines them to gfx1250 builds whose SDK ships the TDM descriptor header.
@@ -250,7 +257,14 @@ ncclResult_t ncclSymkInitOnce(struct ncclComm* comm) {
     struct ncclDevResourceRequirements ginOutboxReq = {};
     struct ncclDevResourceRequirements rsGinAccumReq = {};
     struct ncclDevResourceRequirements railSignalReq = {};
-    if (ncclParamSymGinKernelsEnable() && ncclTeamLsa(comm).nRanks < comm->nRanks) {
+    const bool needGin = ncclTeamLsa(comm).nRanks < comm->nRanks;
+    symk->ginBackendUnsupported = needGin && ncclSymkGinBackendUnsupported(comm);
+    if (symk->ginBackendUnsupported && ncclParamSymGinKernelsEnable()) {
+      INFO(NCCL_INIT, "Symmetric GIN kernels disabled: the Anvil SDMA GIN backend cannot reach rail peers outside "
+                      "the LSA team (lsaSize=%d, nRanks=%d)",
+           ncclTeamLsa(comm).nRanks, comm->nRanks);
+    }
+    if (ncclParamSymGinKernelsEnable() && needGin && !symk->ginBackendUnsupported) {
       int maxBlocks;
       size_t bufSize;
       getRequirements_gin(comm, &maxBlocks, &bufSize);
@@ -394,7 +408,7 @@ uint32_t ncclSymkMask(struct ncclComm* comm, ncclFunc_t coll, int /*ncclDevRedOp
   // Tma kernel degrades to normal selection instead of ending up with an empty mask.
   if (ncclSymkTmaForced(comm) && (kmask & kernelMask_Tma) != 0) kmask &= kernelMask_Tma;
 
-  bool hasGin = ncclParamSymGinKernelsEnable() != 0;
+  bool hasGin = ncclParamSymGinKernelsEnable() != 0 && !comm->symkState.ginBackendUnsupported;
   if (!hasGin) kmask &= ~kernelMask_Gin;
   bool needGin = ncclTeamLsa(comm).nRanks < comm->nRanks;
   kmask &= needGin ? kernelMask_Gin : ~kernelMask_Gin;

projects/rccl/test/gin/GinAnvilPlugin_test.cpp

@@ -90,12 +90,12 @@ class GinAnvilPluginTest : public ::testing::Test {
     ncclGinAnvilSetInitContext(*ictx, mockComm_.get());
   }
 
-  void connectColl(void* ictx, void** coll, int nranks = 1) {
+  void connectColl(void* ictx, void** coll, int nranks = 1, int ginRank = 0) {
     void* listen = nullptr;
     char handle[NCCL_NET_HANDLE_MAXSIZE] = {};
     ASSERT_EQ(plugin_.listen(ictx, 0, handle, &listen), ncclSuccess);
     std::vector<void*> handles(static_cast<size_t>(nranks), handle);
-    ASSERT_EQ(plugin_.connect(ictx, handles.data(), nranks, 0, listen, coll), ncclSuccess);
+    ASSERT_EQ(plugin_.connect(ictx, handles.data(), nranks, ginRank, listen, coll), ncclSuccess);
     ASSERT_EQ(plugin_.closeListen(listen), ncclSuccess);
   }
 
@@ -125,6 +125,27 @@ class GinAnvilPluginTest : public ::testing::Test {
     if (coll) plugin_.closeColl(coll);
     if (ictx) plugin_.finalize(ictx);
   }
+
+  // A refused bind must leave nothing behind: no signals on the device context,
+  // no IPC table entry, and no conn-check traffic.
+  void expectSignalsUnbound(ncclNetDeviceHandle_v11_t* devHandle) {
+    const size_t nContexts = devHandle->size / sizeof(ncclGinAnvilSdmaGPUContext);
+    ASSERT_GE(nContexts, 1u);
+    std::vector<ncclGinAnvilSdmaGPUContext> hostCtx(nContexts);
+    ASSERT_EQ(hipMemcpy(hostCtx.data(), devHandle->handle, nContexts * sizeof(ncclGinAnvilSdmaGPUContext),
+                        hipMemcpyDeviceToHost),
+              hipSuccess);
+    for (size_t c = 0; c < nContexts; c++) {
+      EXPECT_EQ(hostCtx[c].signals, nullptr) << "context " << c;
+      EXPECT_EQ(hostCtx[c].signal_remote_addrs, nullptr) << "context " << c;
+    }
+    const ncclGinAnvilIpcBufEntry* table = nullptr;
+    int ipcCount = -1;
+    ncclGinAnvilIpcTableGetDevice(&table, &ipcCount);
+    EXPECT_EQ(ipcCount, 0);
+    EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckWriteCalls(), 0);
+    EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckVerifyCalls(), 0);
+  }
 };
 
 // G1: wrong NCCL_GIN_TYPE.
@@ -722,8 +743,9 @@ TEST_F(GinAnvilPluginTest, CloseColl_AfterSignalBind) {
   plugin_.finalize(ictx);
 }
 
-// G20: conn-check skipped when nRanks != lsaSize (no gate on single-rank LSA team).
-TEST_F(GinAnvilPluginTest, ConnCheck_SkippedWhenLsaSizeMismatch) {
+// G20: a GIN team wider than the LSA team fails registration. On the first rank of
+// an LSA team, GIN peer 1 would point past the end of a one-slot LSA window.
+TEST_F(GinAnvilPluginTest, BindSignals_GinPeerPastLsaWindowFailsRegistration) {
   GinAnvilPluginStubs::SetBootstrapNranks(2);
   mockComm_.get()->devrState.lsaSize = 1;
   void* ictx = nullptr;
@@ -737,9 +759,38 @@ TEST_F(GinAnvilPluginTest, ConnCheck_SkippedWhenLsaSizeMismatch) {
   ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess);
 
   char arena[4096] = {};
-  EXPECT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 1, 2), ncclSuccess);
-  EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckWriteCalls(), 0);
-  EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckVerifyCalls(), 0);
+  EXPECT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 1, 2), ncclInvalidUsage);
+  expectSignalsUnbound(devHandle);
+
+  plugin_.destroyContext(ginCtx);
+  plugin_.closeColl(coll);
+  plugin_.finalize(ictx);
+}
+
+// The MI300X CPX shape from AICOMRCCL-2706, scaled down: four ranks in two LSA teams
+// of two, seen from world rank 2 (lsaSelf 0). GIN peers 0 and 1 would point before
+// the start of this rank's LSA window.
+TEST_F(GinAnvilPluginTest, BindSignals_GinPeerBeforeLsaWindowFailsRegistration) {
+  GinAnvilPluginStubs::SetBootstrapNranks(4);
+  mockComm_.get()->rank = 2;
+  mockComm_.get()->nRanks = 4;
+  mockComm_.get()->devrState.lsaSelf = 0;
+  mockComm_.get()->devrState.lsaSize = 2;
+  mockComm_.lsaRanks[0] = 2;
+  mockComm_.lsaRanks[1] = 3;
+  void* ictx = nullptr;
+  initCtx(&ictx);
+  void* coll = nullptr;
+  connectColl(ictx, &coll, 4, 2);
+  ncclGinConfig_t cfg{};
+  cfg.nSignals = 4;
+  void* ginCtx = nullptr;
+  ncclNetDeviceHandle_v11_t* devHandle = nullptr;
+  ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess);
+
+  char arena[4096] = {};
+  EXPECT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 1, 4), ncclInvalidUsage);
+  expectSignalsUnbound(devHandle);
 
   plugin_.destroyContext(ginCtx);
   plugin_.closeColl(coll);
@@ -771,14 +822,15 @@ TEST_F(GinAnvilPluginTest, ConnCheck_SkippedWhenSignalSlotTooSmall) {
   plugin_.finalize(ictx);
 }
 
-// Conn-check is an LSA cuMem-VMM probe. Under NCCL_GIN_CONNECTION_RAIL the GIN team
-// indexes rail ranks, not lsaSelf, so skip even when nRanks == lsaSize.
-TEST_F(GinAnvilPluginTest, ConnCheck_SkippedOnRailConnection) {
+// Under NCCL_GIN_CONNECTION_RAIL the GIN peers are ranks on other hosts. Two hosts of
+// two ranks, seen from world rank 0: rail peer 1 is world rank 2, but LSA slot 1 is
+// world rank 1, so registration fails even though nRanks == lsaSize.
+TEST_F(GinAnvilPluginTest, BindSignals_RailGinTeamFailsRegistration) {
   GinAnvilPluginStubs::SetBootstrapNranks(2);
   mockComm_.get()->globalGinSupport = NCCL_GIN_CONNECTION_RAIL;
+  mockComm_.get()->contiguousRanksPerHost = 2;
+  mockComm_.get()->nRanks = 4;
   mockComm_.get()->devrState.lsaSize = 2;
-  mockComm_.get()->devrState.lsaSelf = 0;
-  mockComm_.get()->rank = 1;
   void* ictx = nullptr;
   initCtx(&ictx);
   void* coll = nullptr;
@@ -790,9 +842,8 @@ TEST_F(GinAnvilPluginTest, ConnCheck_SkippedOnRailConnection) {
   ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess);
 
   char arena[4096] = {};
-  EXPECT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 1, 2), ncclSuccess);
-  EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckWriteCalls(), 0);
-  EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckVerifyCalls(), 0);
+  EXPECT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 1, 2), ncclInvalidUsage);
+  expectSignalsUnbound(devHandle);
 
   plugin_.destroyContext(ginCtx);
   plugin_.closeColl(coll);
@@ -801,19 +852,20 @@ TEST_F(GinAnvilPluginTest, ConnCheck_SkippedOnRailConnection) {
 
 // The span-base probe allgathers over comm->bootstrap, which writes one slot per
 // world rank. Under NCCL_GIN_CONNECTION_RAIL the world is wider than the GIN team,
-// so a buffer sized by the GIN team would be overrun.
+// so a buffer sized by the GIN team would be overrun. A wider rail team would fail
+// registration, so this uses one host of 8 ranks, where the rail team is this rank.
 TEST_F(GinAnvilPluginTest, BindSignals_SpanBaseProbeSizedForWorldRanks) {
   constexpr int kWorldRanks = 8;
-  GinAnvilPluginStubs::SetBootstrapNranks(2);
   mockComm_.get()->globalGinSupport = NCCL_GIN_CONNECTION_RAIL;
+  mockComm_.get()->contiguousRanksPerHost = kWorldRanks;
   mockComm_.get()->nRanks = kWorldRanks;
   mockComm_.get()->rank = 5;
-  mockComm_.get()->devrState.lsaSize = 2;
-  mockComm_.get()->devrState.lsaSelf = 0;
+  mockComm_.get()->devrState.lsaSize = kWorldRanks;
+  mockComm_.get()->devrState.lsaSelf = 5;
   void* ictx = nullptr;
   initCtx(&ictx);
   void* coll = nullptr;
-  connectColl(ictx, &coll, 2);
+  connectColl(ictx, &coll);
   ncclGinConfig_t cfg{};
   cfg.nSignals = 2;
   void* ginCtx = nullptr;
@@ -1327,22 +1379,27 @@ TEST_F(GinAnvilPluginTest, ConnCheck_UsesLsaTeamForBootstrap) {
   ASSERT_EQ(hipMalloc(&rawDevLsa, sizeof(uint64_t) * 2), hipSuccess);
   HipAllocation devLsa(rawDevLsa);
   GinAnvilPluginStubs::SetLsaSelfAddr(devLsa.get());
+  GinAnvilPluginStubs::SetBootstrapNranks(2);
+  mockComm_.get()->rank = 1;
   mockComm_.get()->devrState.lsaSelf = 1;
-  mockComm_.lsaRanks[0] = 1;
-  mockComm_.lsaRanks[1] = 0;
 
   void* ictx = nullptr;
+  initCtx(&ictx);
   void* coll = nullptr;
+  connectColl(ictx, &coll, 2, 1);
+  ncclGinConfig_t cfg{};
+  cfg.nSignals = 2;
   void* ginCtx = nullptr;
-  startTwoRankGin(&ictx, &coll, &ginCtx);
+  ncclNetDeviceHandle_v11_t* devHandle = nullptr;
+  ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess);
 
   char arena[4096] = {};
   EXPECT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 1, 2), ncclSuccess);
   EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeAllGatherRank(), 1);
   EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeAllGatherNranks(), 2);
-  EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeAllGatherRanks(), std::vector<int>({1, 0}));
+  EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeAllGatherRanks(), std::vector<int>({0, 1}));
   EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeBarrierRank(), 1);
-  EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeBarrierRanks(), std::vector<int>({1, 0}));
+  EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeBarrierRanks(), std::vector<int>({0, 1}));
 
   stopGin(ictx, coll, ginCtx);
 }

projects/rccl/test/host/sym-kernels-test.cc

@@ -36,6 +36,13 @@
 
 #include "SymTdmTestHelpers.h"
 
+// sym_kernels.cc asks for the railed GIN backend; gin_host.cc is not linked into this binary.
+std::function<ncclGinType_t(ncclComm*)> g_railedGinType = [](ncclComm*) { return NCCL_GIN_TYPE_NONE; };
+ncclResult_t ncclGetRailedGinType(struct ncclComm* comm, ncclGinType_t* ginType) {
+  *ginType = g_railedGinType(comm);
+  return ncclSuccess;
+}
+
 namespace {
 
 bool KernelBitSet(uint32_t mask, ncclSymkKernelId id) { return (mask >> static_cast<int>(id)) & 1u; }
@@ -369,6 +376,27 @@ TEST_F(SymKernelMaskTest, NeedGinTrue_HasGinFalse_EverythingCleared) {
   EXPECT_EQ(kmask, 0u);
 }
 
+// Init marks the GIN backend unsupported (Anvil SDMA): the same empty mask as SYM_GIN_KERNELS_ENABLE=0.
+TEST_F(SymKernelMaskTest, NeedGinTrue_GinBackendUnsupported_EverythingCleared) {
+  ScopedHook teamLsa(g_ncclTeamLsa, [](ncclComm_t) {
+    ncclTeam_t t{};
+    t.nRanks = 2;
+    return t;
+  });
+  ScopedHook loadParam(g_loadParam, SymGinKernelsEnable(1));
+  comm_->symkState.ginBackendUnsupported = true;
+  uint32_t kmask = ncclSymkMask(comm_.get(), ncclFuncReduceScatter, ncclDevSum, kTy, /*nElts=*/1024);
+  EXPECT_EQ(kmask, 0u);
+}
+
+TEST_F(SymKernelMicrotest, GinBackendUnsupported_OnlyForAnvilSdma) {
+  for (ncclGinType_t type : {NCCL_GIN_TYPE_ANVIL_SDMA, NCCL_GIN_TYPE_PROXY, NCCL_GIN_TYPE_NONE}) {
+    ScopedHook railed(g_railedGinType, [type](ncclComm*) { return type; });
+    EXPECT_EQ(ncclSymkGinBackendUnsupported(comm_.get()), type == NCCL_GIN_TYPE_ANVIL_SDMA) << "ginType " << type;
+    EXPECT_EQ(railed.calls, 1);
+  }
+}
+
 TEST_F(SymKernelMaskTest, HasLsaMultimemTrue_STMCKernelSurvives) {
   comm_->symkState.hasLsaMultimem = true;
   uint32_t kmask = ncclSymkMask(comm_.get(), ncclFuncAllReduce, ncclDevSum, kTy, /*nElts=*/1024);