dlamd1dai · Oct 31, 2026 · Oct 24, 2026 · Oct 25, 2026 · Oct 27, 2026 · Oct 27, 2026
@@ -175,6 +175,16 @@ static bool ginGateAllows(uint64_t ginTypeBitMask, const struct ncclComm* comm)
   return comm->nNodes <= 1;
 }
 
+// GIN connection type once ncclGinSetDefaultBackend has run. That can drop every active backend
+// (Anvil SDMA alone on a multi-host communicator), and FULL or RAIL with none left would make
+// query and connect report a backend that is gone.
+static ncclGinConnectionType_t ginConnectionType(const struct ncclComm* comm, bool globalCrossNicSupport) {
+  if (!comm->sharedRes->ginState.supported) return NCCL_GIN_CONNECTION_NONE;
+  if (globalCrossNicSupport) return NCCL_GIN_CONNECTION_FULL;
+  if (comm->contiguousRanksPerHost != INT_MAX) return NCCL_GIN_CONNECTION_RAIL;
+  return NCCL_GIN_CONNECTION_NONE;
+}
+
 static bool ctaPolicyIsValid(int ctaPolicy) {
   int availCtaPolicies[3] = {NCCL_CTA_POLICY_DEFAULT, NCCL_CTA_POLICY_EFFICIENCY, NCCL_CTA_POLICY_ZERO};
   int maxPolicy = 0;
@@ -2863,11 +2873,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
   comm->globalGinSupport = NCCL_GIN_CONNECTION_NONE;
   if (ginGateAllows(globalGinTypeBitMask, comm)) {
     NCCLCHECKGOTO(ncclGinSetDefaultBackend(comm, globalGinTypeBitMask), ret, fail);
-    if (globalCrossNicSupport) {
-      comm->globalGinSupport = NCCL_GIN_CONNECTION_FULL;
-    } else if (comm->contiguousRanksPerHost != INT_MAX) {
-      comm->globalGinSupport = NCCL_GIN_CONNECTION_RAIL;
-    }
+    comm->globalGinSupport = ginConnectionType(comm, globalCrossNicSupport);
   }
   comm->globalRmaProxySupport = globalRmaPluginSupport && globalCrossNicSupport && comm->cuMemGdrSupport;
   isOneLsaTeams = ncclDevrIsOneLsaTeam(comm);

@@ -12,6 +12,7 @@
 #include "checks.h"
 #include "plugin.h"
 #include "nccl_gin.h"
+#include "graph/topo.h"
 #include "gin/gin_host_proxy.h"
 #ifdef ENABLE_ROCSHMEM_GIN
 #include "gin/gin_host_rocshmem_gda.h"
@@ -446,11 +447,21 @@ ncclResult_t ncclGinSetDefaultBackend(struct ncclComm* comm, uint64_t globalBitM
   if (ginState->finalized) return ncclSuccess;
   std::lock_guard<std::mutex> lock(pluginMutex);
 
-  // Close ALL non-global backends and compact the array
+  // Close ALL non-global backends and compact the array. Anvil SDMA addresses
+  // peers by local device id, so it is dropped when more than one physical host
+  // is in the communicator (including multi-host MNNVL where nNodes folds to 1).
   for (int i = 0; i < initialNumActiveBackends; i++) {
     struct ncclGinBackendState* backend = &ginState->backends[i];
-    if (!(BIT(backend->ginType) & globalBitMask)) {
-      WARN("GIN backend type %d not supported on all ranks, removing from active backends", backend->ginType);
+    const bool anvilMultiHost =
+        comm->topo->nHosts > 1 && backend->ginType == NCCL_GIN_TYPE_ANVIL_SDMA;
+    if (anvilMultiHost || !(BIT(backend->ginType) & globalBitMask)) {
+      if (anvilMultiHost) {
+        INFO(NCCL_INIT | NCCL_NET,
+             "GIN backend ANVIL_SDMA is single-host only; removing it from active backends (nHosts=%d)",
+             comm->topo->nHosts);
+      } else {
+        WARN("GIN backend type %d not supported on all ranks, removing from active backends", backend->ginType);
+      }
       NCCLCHECK(ncclGinPluginFinalize(comm, backend->pluginIndex, backend->ginInstance));
     } else {
       ginState->backends[currentNumActiveBackends++] = *backend;

@@ -26,6 +26,7 @@
 
 #include "nccl.h"
 #include "comm.h"
+#include "graph/topo.h"
 #include "nccl_gin.h"
 #include "plugin.h"
 
@@ -166,13 +167,16 @@ TEST_F(GinPluginInitTest, DoesNotFinalizeWhenInitPointerIsNull) {
 class GinFinalizeTest : public GinPluginInitTest {
  protected:
   std::unique_ptr<ncclSharedResources> sharedRes_ = std::make_unique<ncclSharedResources>();
+  std::unique_ptr<ncclTopoSystem> topo_ = std::make_unique<ncclTopoSystem>();
 
   void SetUp() override {
     GinPluginInitTest::SetUp();
     pluginLib_.state = ncclGinPluginStateEnabled;
     pluginLibs[0] = pluginLib_;
     pluginLibs[0].refCount = kInitialRefCount;
     comm_.sharedRes = sharedRes_.get();
+    topo_->nHosts = 1;
+    comm_.topo = topo_.get();
   }
 
   void TearDown() override {
@@ -247,6 +251,69 @@ TEST_F(GinFinalizeTest, LiveBackend_SetDefaultBackendArmsSupported) {
   EXPECT_EQ(fake_.finalizeCalls, 0);
 }
 
+// AICOMRCCL-3417: Anvil SDMA stays on a single physical host alongside PROXY.
+TEST_F(GinFinalizeTest, AnvilSdma_KeptWithProxyOnSingleHost) {
+  void* anvil = std::malloc(8);
+  void* proxy = std::malloc(8);
+  fake_.lastCtx = anvil;
+  addActiveBackend(anvil);
+  addActiveBackend(proxy);
+  struct ncclGinState* ginState = &sharedRes_->ginState;
+  ginState->backends[0].ginType = NCCL_GIN_TYPE_ANVIL_SDMA;
+  ginState->backends[1].ginType = NCCL_GIN_TYPE_PROXY;
+  topo_->nHosts = 1;
+
+  EXPECT_EQ(ncclGinSetDefaultBackend(&comm_,
+                                     BIT(NCCL_GIN_TYPE_ANVIL_SDMA) | BIT(NCCL_GIN_TYPE_PROXY)),
+            ncclSuccess);
+  EXPECT_TRUE(ginState->supported);
+  EXPECT_EQ(ginState->numActiveBackends, 2);
+  EXPECT_EQ(fake_.finalizeCalls, 0);
+  std::free(anvil);
+  std::free(proxy);
+  fake_.lastCtx = nullptr;
+}
+
+// AICOMRCCL-3417: Anvil SDMA is dropped when topo reports more than one host,
+// including MNNVL pods that fold nNodes to 1. PROXY remains the default.
+TEST_F(GinFinalizeTest, AnvilSdma_DroppedOnMultiHostKeepsProxy) {
+  void* anvil = std::malloc(8);
+  void* proxy = std::malloc(8);
+  fake_.lastCtx = anvil;
+  addActiveBackend(anvil);
+  addActiveBackend(proxy);
+  struct ncclGinState* ginState = &sharedRes_->ginState;
+  ginState->backends[0].ginType = NCCL_GIN_TYPE_ANVIL_SDMA;
+  ginState->backends[1].ginType = NCCL_GIN_TYPE_PROXY;
+  topo_->nHosts = 2;
+
+  EXPECT_EQ(ncclGinSetDefaultBackend(&comm_,
+                                     BIT(NCCL_GIN_TYPE_ANVIL_SDMA) | BIT(NCCL_GIN_TYPE_PROXY)),
+            ncclSuccess);
+  EXPECT_TRUE(ginState->supported);
+  EXPECT_EQ(ginState->numActiveBackends, 1);
+  EXPECT_EQ(ginState->backends[0].ginType, NCCL_GIN_TYPE_PROXY);
+  EXPECT_EQ(fake_.finalizeCalls, 1);
+  std::free(proxy);
+  fake_.lastCtx = nullptr;
+}
+
+// AICOMRCCL-3417: Anvil-only on multi-host leaves zero backends and clears supported
+// so init.cc keeps globalGinSupport NONE.
+TEST_F(GinFinalizeTest, AnvilSdmaOnly_MultiHostClearsSupported) {
+  void* anvil = std::malloc(8);
+  fake_.lastCtx = anvil;
+  addActiveBackend(anvil);
+  struct ncclGinState* ginState = &sharedRes_->ginState;
+  ginState->backends[0].ginType = NCCL_GIN_TYPE_ANVIL_SDMA;
+  topo_->nHosts = 2;
+
+  EXPECT_EQ(ncclGinSetDefaultBackend(&comm_, BIT(NCCL_GIN_TYPE_ANVIL_SDMA)), ncclSuccess);
+  EXPECT_FALSE(ginState->supported);
+  EXPECT_EQ(ginState->numActiveBackends, 0);
+  EXPECT_EQ(fake_.finalizeCalls, 1);
+}
+
 // Success path: keep the context, mark the plugin enabled, do not finalize.
 TEST_F(GinPluginInitTest, KeepsContextWhenDevicesSucceed) {
   fake_.ndev = 2;

@@ -3095,6 +3095,30 @@ TEST_F(InitMicrotest, GinGate_NoSupportedGinTypeOrNoCuMemGdr_RefusesGin) {
   EXPECT_FALSE(ginGateAllows(UINT64_MAX, noGdr.get()));
 }
 
+// The connection type chosen after ncclGinSetDefaultBackend. AICOMRCCL-3417: dropping Anvil SDMA on
+// a multi-host communicator can leave no backend, and that communicator must stay NONE.
+TEST_F(InitMicrotest, GinConnectionType_NoBackendLeft_StaysNone) {
+  auto comm = std::make_unique<ncclComm>();
+  auto sr = std::make_unique<ncclSharedResources>();
+  comm->sharedRes = sr.get();
+  comm->contiguousRanksPerHost = 8;
+  sr->ginState.supported = false;
+  EXPECT_EQ(ginConnectionType(comm.get(), /*globalCrossNicSupport=*/true), NCCL_GIN_CONNECTION_NONE);
+  EXPECT_EQ(ginConnectionType(comm.get(), /*globalCrossNicSupport=*/false), NCCL_GIN_CONNECTION_NONE);
+}
+
+TEST_F(InitMicrotest, GinConnectionType_BackendLeft_FullOrRail) {
+  auto comm = std::make_unique<ncclComm>();
+  auto sr = std::make_unique<ncclSharedResources>();
+  comm->sharedRes = sr.get();
+  sr->ginState.supported = true;
+  comm->contiguousRanksPerHost = 8;
+  EXPECT_EQ(ginConnectionType(comm.get(), /*globalCrossNicSupport=*/true), NCCL_GIN_CONNECTION_FULL);
+  EXPECT_EQ(ginConnectionType(comm.get(), /*globalCrossNicSupport=*/false), NCCL_GIN_CONNECTION_RAIL);
+  comm->contiguousRanksPerHost = INT_MAX;
+  EXPECT_EQ(ginConnectionType(comm.get(), /*globalCrossNicSupport=*/false), NCCL_GIN_CONNECTION_NONE);
+}
+
 // NOT ASSERTABLE FROM THIS RUNG, deliberately: the four `global*Support` accumulators at :1491-1494 are
 // function-locals first read at :2347-2363, ~700 lines past the terminator. They execute (so they count
 // as covered) but nothing here can observe them, and deleting any of the four leaves the suite green.

@@ -19,10 +19,12 @@
 #include "SymmetricMemPrereq.hpp"
 #include "TestChecks.hpp"
 #include "comm.h"
+#include "graph/topo.h"
 
 #include "nccl_device.h"
 #include "nccl_device/gin/anvil_sdma/gin_anvil_sdma_device_host_common.h"
 #ifdef ENABLE_ROCSHMEM_GIN
+#include "gin/gin_anvil_sdma_factory.h"
 #include "gin/gin_host_anvil_sdma.h"
 #endif
 
@@ -135,6 +137,29 @@ std::string anvilSingleNodeReason() {
   return "";
 }
 
+// AICOMRCCL-3417. The rest of this file requires NCCL_GIN_TYPE. This case is
+// the opposite: unset, so Anvil SDMA and PROXY are both eligible and SDMA
+// would otherwise sort ahead of PROXY on a multi-host communicator.
+std::string mixedBackendTwoNodeSkipReason() {
+  std::string local = ginEnvDisabledReason();
+  if (local.empty()) local = cuMemReason();
+#ifndef ENABLE_ROCSHMEM_GIN
+  if (local.empty()) local = "AICOMRCCL-3417 requires ENABLE_ROCSHMEM_GIN";
+#else
+  if (local.empty() && gin_anvil_sdma_probe() <= 0)
+    local = "AICOMRCCL-3417 requires Anvil SDMA (gin_anvil_sdma_probe)";
+#endif
+  if (local.empty()) {
+    const char* ginType = std::getenv("NCCL_GIN_TYPE");
+    if (ginType && ginType[0])
+      local = "AICOMRCCL-3417 requires NCCL_GIN_TYPE to be unset";
+  }
+  // crossNodeReason splits MPI_COMM_TYPE_SHARED. Every rank must call it.
+  std::string cross = crossNodeReason();
+  if (local.empty()) local = std::move(cross);
+  return mpiCoordinatedSkipReason(!local.empty(), local.empty() ? nullptr : local.c_str());
+}
+
 // Unlike anvilSingleNodeReason() above, this applies on every backend: the
 // fence-visibility tests pair each rank with one peer and read back the peer's
 // write, so they are only meaningful with all ranks on one node.
@@ -436,7 +461,8 @@ class GinMPIDeviceTests : public MPITestBase {
   // Minimal 64-byte put + waitSignal round-trip from rank 0 to rank 1 on the
   // caller's comm and stream. Contains ASSERT_MPI_*, so it must be the
   // caller's last statement.
-  void runPutWaitRoundTrip(ncclComm_t comm, hipStream_t stream) {
+  void runPutWaitRoundTrip(ncclComm_t comm, hipStream_t stream,
+                           ncclGinType_t ginType = NCCL_GIN_TYPE_NONE) {
     int rank = -1, nRanks = -1;
     ncclCommUserRank(comm, &rank);
     ncclCommCount(comm, &nRanks);
@@ -474,6 +500,7 @@ class GinMPIDeviceTests : public MPITestBase {
     ncclDevCommRequirements reqs = defaultGinReqs();
     reqs.railGinBarrierCount = 1;
     reqs.ginSignalCount      = 1;
+    reqs.ginType             = ginType;
     ncclDevComm devComm{};
     ASSERT_MPI_EQ(ncclSuccess, ncclDevCommCreate(comm, &reqs, &devComm));
     auto devCommCleanup = makeScopeGuard([&]() {
@@ -7867,4 +7894,47 @@ TEST_F(GinProxyNthreadsSplitTests, SplitShare_SecondDevCommPut) {
   runPutWaitRoundTrip(child, getActiveStream());
 }
 
+// Two ranks, two physical hosts, NCCL_GIN_TYPE unset (AICOMRCCL-3417). Anvil
+// SDMA must not stay active: its factory addresses peers by local device id,
+// including across hosts in one MNNVL pod where nNodes folds to 1. PROXY is
+// the default, a symmetric window registers, a default DevComm and an explicit
+// PROXY DevComm each complete a cross-host put/signal, and an explicit Anvil
+// DevComm is rejected.
+TEST_F(GinMPIDeviceTests, MixedBackends_TwoNodeUnsetType) {
+  if (auto reason = mixedBackendTwoNodeSkipReason(); !reason.empty()) GTEST_SKIP() << reason;
+  if (auto reason = mpiCoordinatedSkipReason(
+          !validateTestPrerequisites(/*min_processes=*/2, /*max_processes=*/2,
+                                     /*require_power_of_two=*/false, /*min_nodes=*/2, /*max_nodes=*/2),
+          "Requires exactly 2 ranks on 2 hosts");
+      !reason.empty())
+    GTEST_SKIP() << reason;
+
+  ASSERT_EQ(ncclSuccess, createTestCommunicator());
+  SKIP_IF_GIN_UNSUPPORTED();
+  ncclComm_t comm = getActiveCommunicator();
+  if (auto reason = mpiCoordinatedSkipReason(comm->topo->nHosts < 2, "Requires topo->nHosts >= 2");
+      !reason.empty())
+    GTEST_SKIP() << reason;
+  ASSERT_MPI_EQ(true, comm->topo->nHosts >= 2);
+
+  ncclCommProperties_t props = NCCL_COMM_PROPERTIES_INITIALIZER;
+  ASSERT_MPI_EQ(ncclSuccess, ncclCommQueryProperties(comm, &props));
+  ASSERT_MPI_EQ(NCCL_GIN_TYPE_PROXY, props.ginType);
+  ASSERT_MPI_EQ(true, props.ginSupport[NCCL_GIN_TYPE_PROXY]);
+  ASSERT_MPI_EQ(false, props.ginSupport[NCCL_GIN_TYPE_ANVIL_SDMA]);
+
+  ncclDevCommRequirements anvilReqs = defaultGinReqs();
+  anvilReqs.ginType = NCCL_GIN_TYPE_ANVIL_SDMA;
+  anvilReqs.railGinBarrierCount = 1;
+  anvilReqs.ginSignalCount = 1;
+  ncclDevComm anvil{};
+  ncclResult_t anvilSt = ncclDevCommCreate(comm, &anvilReqs, &anvil);
+  if (anvilSt == ncclSuccess) (void)ncclDevCommDestroy(comm, &anvil);
+  ASSERT_MPI_EQ(ncclInternalError, anvilSt);
+
+  // Default backend (ginType unset on the DevComm) and an explicit PROXY one.
+  runPutWaitRoundTrip(comm, getActiveStream());
+  runPutWaitRoundTrip(comm, getActiveStream(), NCCL_GIN_TYPE_PROXY);
+}
+
 #endif  // MPI_TESTS_ENABLED