From 9aed364cadf99f25d014b07885b3e5aab8b9b00a Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Tue, 13 Oct 2026 21:28:57 -0400 Subject: [PATCH 1/3] fix(rccl/gin): reject Anvil GIN peers outside the LSA team ginAnvilRegisterLsaSignals builds the signal peer table as lsaSelf + (pe - ginRank) * bigSize for every GIN peer. The table is indexed by GIN peer, but each entry is a slot in this rank's LSA window, so it is only valid for peers that are members of the LSA team. When the GIN team is wider than the LSA team (MI300X CPX with 8 ranks and NCCL_LSA_TEAM_SIZE=4, or any RAIL team with more than one rank), entries for the other peers point past either end of the mapped window. The plugin warned and continued, skipped the connectivity check, and left the device to signal through those addresses. Check every GIN peer against lsaRankList before binding, using the RAIL stride for the GIN-to-world mapping, and fail registration with ncclInvalidUsage naming the first peer outside the team. JIRA ID: AICOMRCCL-2706 Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 1 + .../rccl/src/gin/gin_plugin_anvil_sdma.cc | 30 +++++-- .../rccl/test/gin/GinAnvilPlugin_test.cpp | 90 ++++++++++++++----- 3 files changed, 95 insertions(+), 26 deletions(-) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index ae09dc8b1a77..b7daac6abfaf 100644 --- a/projects/rccl/CHANGELOG.md +++ b/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. When the LSA team is smaller than the communicator, the symmetric kernels bind GIN signals during communicator initialization, so `ncclCommInitRank` fails; with `NCCL_SYM_GIN_KERNELS_ENABLE=0`, `ncclDevCommCreate` fails instead. These configurations previously completed initialization; set `NCCL_LSA_TEAM_SIZE` to cover the GIN team or select another GIN type. ### Known issues diff --git a/projects/rccl/src/gin/gin_plugin_anvil_sdma.cc b/projects/rccl/src/gin/gin_plugin_anvil_sdma.cc index 1cf2dfed7aff..c1b816ff2be2 100644 --- a/projects/rccl/src/gin/gin_plugin_anvil_sdma.cc +++ b/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); diff --git a/projects/rccl/test/gin/GinAnvilPlugin_test.cpp b/projects/rccl/test/gin/GinAnvilPlugin_test.cpp index 5c0da9587623..3ed0a455ce5f 100644 --- a/projects/rccl/test/gin/GinAnvilPlugin_test.cpp +++ b/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 handles(static_cast(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,21 @@ 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) { + ncclGinAnvilSdmaGPUContext hostCtx{}; + ASSERT_EQ(hipMemcpy(&hostCtx, devHandle->handle, sizeof(hostCtx), hipMemcpyDeviceToHost), hipSuccess); + EXPECT_EQ(hostCtx.signals, nullptr); + EXPECT_EQ(hostCtx.signal_remote_addrs, nullptr); + 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 +737,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 +753,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 +816,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 +836,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); @@ -1327,22 +1372,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({1, 0})); + EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeAllGatherRanks(), std::vector({0, 1})); EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeBarrierRank(), 1); - EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeBarrierRanks(), std::vector({1, 0})); + EXPECT_EQ(GinAnvilPluginStubs::GetLastIntraNodeBarrierRanks(), std::vector({0, 1})); stopGin(ictx, coll, ginCtx); } From 9b3051c71eacec8c5b2aef3a81b1e9b60186c25e Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Mon, 26 Oct 2026 16:22:43 -0400 Subject: [PATCH 2/3] test(rccl/gin): adapt Anvil LSA team check tests to per-context signal spans The check now runs in ginAnvilRegisterSignalSpan, once per resource-window bind. expectSignalsUnbound checks every logical context of the handle, and BindSignals_SpanBaseProbeSizedForWorldRanks uses a one-host rail team so it still passes registration while the world stays wider than the GIN team. Co-authored-by: Cursor --- .../rccl/test/gin/GinAnvilPlugin_test.cpp | 25 ++++++++++++------- 1 file changed, 16 insertions(+), 9 deletions(-) diff --git a/projects/rccl/test/gin/GinAnvilPlugin_test.cpp b/projects/rccl/test/gin/GinAnvilPlugin_test.cpp index 3ed0a455ce5f..225ce7a3eda3 100644 --- a/projects/rccl/test/gin/GinAnvilPlugin_test.cpp +++ b/projects/rccl/test/gin/GinAnvilPlugin_test.cpp @@ -129,10 +129,16 @@ class GinAnvilPluginTest : public ::testing::Test { // 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) { - ncclGinAnvilSdmaGPUContext hostCtx{}; - ASSERT_EQ(hipMemcpy(&hostCtx, devHandle->handle, sizeof(hostCtx), hipMemcpyDeviceToHost), hipSuccess); - EXPECT_EQ(hostCtx.signals, nullptr); - EXPECT_EQ(hostCtx.signal_remote_addrs, nullptr); + const size_t nContexts = devHandle->size / sizeof(ncclGinAnvilSdmaGPUContext); + ASSERT_GE(nContexts, 1u); + std::vector 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); @@ -846,19 +852,20 @@ TEST_F(GinAnvilPluginTest, BindSignals_RailGinTeamFailsRegistration) { // 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; From 126528dfdfe77e3f2d7522628ba215f8b6bedc9f Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Tue, 27 Oct 2026 18:02:57 -0400 Subject: [PATCH 3/3] fix(rccl/symk): skip symmetric GIN kernels on the Anvil SDMA backend The symmetric GIN kernels address the rail team, whose peers are all outside the LSA team. Anvil SDMA reaches a peer only through its slot in the local LSA window, so it can never serve them, and requesting them made Anvil's LSA team check fail ncclCommInitRank for applications that never use GIN. Init now records an unsupported GIN backend and leaves the GIN requirement and kernels out, as NCCL_SYM_GIN_KERNELS_ENABLE=0 does. Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 2 +- projects/rccl/src/include/sym_kernels.h | 1 + projects/rccl/src/sym_kernels.cc | 18 +++++++++++-- projects/rccl/test/host/sym-kernels-test.cc | 28 +++++++++++++++++++++ 4 files changed, 46 insertions(+), 3 deletions(-) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index b7daac6abfaf..b004e206c8e0 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -15,7 +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. When the LSA team is smaller than the communicator, the symmetric kernels bind GIN signals during communicator initialization, so `ncclCommInitRank` fails; with `NCCL_SYM_GIN_KERNELS_ENABLE=0`, `ncclDevCommCreate` fails instead. These configurations previously completed initialization; set `NCCL_LSA_TEAM_SIZE` to cover the GIN team or select another GIN type. +* 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 diff --git a/projects/rccl/src/include/sym_kernels.h b/projects/rccl/src/include/sym_kernels.h index 1c9eb48b7123..8a20c84ef84a 100644 --- a/projects/rccl/src/include/sym_kernels.h +++ b/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; }; diff --git a/projects/rccl/src/sym_kernels.cc b/projects/rccl/src/sym_kernels.cc index 075aaf41ebb9..128a9031f948 100644 --- a/projects/rccl/src/sym_kernels.cc +++ b/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; diff --git a/projects/rccl/test/host/sym-kernels-test.cc b/projects/rccl/test/host/sym-kernels-test.cc index 89f7136be74d..09cdf2991500 100644 --- a/projects/rccl/test/host/sym-kernels-test.cc +++ b/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 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(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);