[RCCL] Enforce model index matching across nodes (#5377)

## Motivation

RCCL relies on a single, consistent Rome / GIO preset topology model
index across all ranks so precomputed graphs (rings, trees, etc.) stay
in lockstep. If different ranks infer different romeTopoModelIdx,
behavior is undefined and jobs should fail fast with actionable logs.

## Technical Details

Use majority-style reference (plurality vote) for deciding reference
index. For mismatched ranks, one line will be printed per physical host.

## JIRA ID

## Test Plan

Manually introduce mismatched nodes and check output.

## Test Result

rocm-systems/projects/rccl/build/hipify/src/init.cc:1681 NCCL WARN RCCL
FATAL: mismatched Rome preset topology model index across ranks; all
ranks must agree for precomputed graphs (voted refIdx 40 from 16 of 24
ranks).

rocm-systems/projects/rccl/build/hipify/src/init.cc:1693 NCCL WARN rank
8 host useocpm2m-097-019 romeTopoModelIdx=-1


## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

---------

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

Author: wenkaidu

projects/rccl/src/CMakeLists.txt

@@ -64,6 +64,8 @@ set(SRC_FILES
   graph/rings.h
   graph/rome_models.cc
   graph/rome_models.h
+  graph/rome_topo_consensus.cc
+  graph/rome_topo_consensus.h
   graph/search.cc
   graph/topo.cc
   graph/topo.h

projects/rccl/src/graph/rome_models.cc

@@ -2080,6 +2080,17 @@ static void parseOptions(struct ncclTopoSystem* system, const char *options) {
   }
 }
 
+static ncclResult_t rcclTopoSetPresetRomeModelIdx(struct ncclTopoSystem* system, int idx) {
+  if (system->romeTopoModelIdx == RCCL_ROME_TOPO_PRESET_MODEL_IDX_NONE) {
+    system->romeTopoModelIdx = idx;
+  } else if (system->romeTopoModelIdx != idx) {
+    WARN("RCCL: conflicting preset Rome topology model index on this node (already %d, attempted %d).",
+         system->romeTopoModelIdx, idx);
+    return ncclInvalidUsage;
+  }
+  return ncclSuccess;
+}
+
 static bool checkOption(const char *options, const char *name) {
   if (strcmp(options, "")) {
     char *str_temp = (char *)malloc(strlen(options) + 1);
@@ -2666,6 +2677,10 @@ ncclResult_t parseA2a8P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
     break;
   }
 
+  if (graph->nChannels) {
+    NCCLCHECK(rcclTopoSetPresetRomeModelIdx(system, i));
+  }
+
   // clean up
   free(all_gpu_permutations);
   return ncclSuccess;
@@ -2833,6 +2848,9 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph*
     applyNetOverride(system, romeTopoModels[i].options);
     break;
   }
+  if (graph->nChannels) {
+    NCCLCHECK(rcclTopoSetPresetRomeModelIdx(system, i));
+  }
   return ncclSuccess;
 }
 
@@ -2972,6 +2990,9 @@ ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
   NCCLCHECK(parseGraph(romeTopoModels[i].ringBase, system, graph, rdm, nnets > 1 ? n : NULL, false));
 
   if (romeTopoModels[i].treeBase != nullptr) NCCLCHECK(parseGraphLight(romeTopoModels[i].treeBase, system, graph, rdm));
+  if (graph->nChannels) {
+    NCCLCHECK(rcclTopoSetPresetRomeModelIdx(system, i));
+  }
   // clean up
   free(all_gpu_permutations);
   return ncclSuccess;
@@ -3092,6 +3113,9 @@ ncclResult_t parse4H4P(struct ncclTopoSystem* system, struct ncclTopoGraph* grap
   parseOptions(system, rome_model_68.options);
   // create 4P4H based on reference and remapped ids
   NCCLCHECK(parseGraph(rome_model_68.ringBase, system, graph, rdm, n_hives.data(), false));
+  if (graph->nChannels) {
+    NCCLCHECK(rcclTopoSetPresetRomeModelIdx(system, RCCL_ROME_TOPO_PRESET_MODEL_IDX_4H4P));
+  }
   return ncclSuccess;
 }
 
@@ -3170,6 +3194,9 @@ ncclResult_t parseGIOTopos(struct ncclTopoSystem* system, struct ncclTopoGraph*
   system->type |= RCCL_TOPO_4P2H_ROME;
 
   NCCLCHECKGOTO(parseGraph(gio16gColumbaModel.ringBase, system, graph, rdm, NULL, false), r, exit);
+  if (graph->nChannels) {
+    NCCLCHECK(rcclTopoSetPresetRomeModelIdx(system, RCCL_ROME_TOPO_PRESET_MODEL_IDX_GIO_COLUMBA));
+  }
 
 exit:
   return ncclSuccess;

projects/rccl/src/graph/rome_topo_consensus.cc

@@ -0,0 +1,69 @@
+/*************************************************************************
+ * Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#include "graph/rome_topo_consensus.h"
+#include "debug.h"
+#include <unordered_map>
+
+#if defined(__GNUC__) || defined(__clang__)
+__attribute__((visibility("default")))
+#endif
+ncclResult_t rcclCheckRomeTopoModelIdxConsensus(
+  int nranks,
+  std::function<int(int)> getRomeTopoModelIdx,
+  std::function<const char*(int)> getHostname,
+  std::function<uint64_t(int)> getHostHash) {
+  if (nranks <= 0) return ncclSuccess;
+
+  std::unordered_map<int, std::pair<int, int>> tallies; // modelIdx -> (vote count, first rank)
+  tallies.reserve(nranks);
+  for (int r = 0; r < nranks; r++) {
+    int v = getRomeTopoModelIdx(r);
+    auto it = tallies.find(v);
+    if (it == tallies.end()) {
+      tallies.emplace(v, std::make_pair(1, r));
+    } else {
+      it->second.first++;
+    }
+  }
+  int refIdx = getRomeTopoModelIdx(0);
+  int refVotes = -1;
+  int refFirstRank = nranks;
+  for (const auto& e : tallies) {
+    int cnt = e.second.first;
+    int firstRank = e.second.second;
+    if (cnt > refVotes || (cnt == refVotes && firstRank < refFirstRank)) {
+      refVotes = cnt;
+      refIdx = e.first;
+      refFirstRank = firstRank;
+    }
+  }
+  if (tallies.size() == 1) return ncclSuccess;
+  int nDisagree = 0;
+  for (int r = 0; r < nranks; r++) {
+    if (getRomeTopoModelIdx(r) != refIdx) nDisagree++;
+  }
+  if (nDisagree > 0) {
+    WARN("RCCL FATAL: mismatched Rome preset topology model index across ranks; all ranks must agree for precomputed graphs (voted refIdx %d from %d of %d ranks).", refIdx, refVotes, nranks);
+    std::unordered_map<uint64_t, int> lowestMismatchRankByHost;
+    lowestMismatchRankByHost.reserve(nranks);
+    for (int r = 0; r < nranks; r++) {
+      if (getRomeTopoModelIdx(r) == refIdx) continue;
+      uint64_t h = getHostHash(r);
+      if (lowestMismatchRankByHost.find(h) == lowestMismatchRankByHost.end()) {
+        lowestMismatchRankByHost.emplace(h, r);
+      }
+    }
+    for (int r = 0; r < nranks; r++) {
+      if (getRomeTopoModelIdx(r) == refIdx) continue;
+      uint64_t h = getHostHash(r);
+      if (lowestMismatchRankByHost[h] != r) continue;
+      WARN("  rank %d host %s romeTopoModelIdx=%d", r, getHostname(r), getRomeTopoModelIdx(r));
+    }
+    return ncclInvalidUsage;
+  }
+  return ncclSuccess;
+}

projects/rccl/src/graph/rome_topo_consensus.h

@@ -0,0 +1,21 @@
+/*************************************************************************
+ * Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#ifndef RCCL_GRAPH_ROME_TOPO_CONSENSUS_H_
+#define RCCL_GRAPH_ROME_TOPO_CONSENSUS_H_
+
+#include "nccl.h"
+#include <functional>
+
+/* Plurality vote on romeTopoModelIdx across ranks; all ranks must match the winner.
+ * Callers supply getters so this stays independent of ncclComm / allGatherInfo. */
+ncclResult_t rcclCheckRomeTopoModelIdxConsensus(
+  int nranks,
+  std::function<int(int /*rank*/)> getRomeTopoModelIdx,
+  std::function<const char*(int /*rank*/)> getHostname,
+  std::function<uint64_t(int /*rank*/)> getHostHash);
+
+#endif

projects/rccl/src/graph/topo.cc

@@ -757,6 +757,7 @@ ncclResult_t ncclTopoAddC2c(struct ncclXmlNode* node, struct ncclTopoSystem* sys
 ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem, const uint64_t localHostHash) {
   NCCLCHECK(ncclCalloc(topoSystem, 1));
   struct ncclTopoSystem* system = *topoSystem;
+  system->romeTopoModelIdx = RCCL_ROME_TOPO_PRESET_MODEL_IDX_NONE;
   struct ncclXmlNode* topNode;
   NCCLCHECK(xmlFindTag(xml, "system", &topNode));
   for (int s=0; s<topNode->nSubs; s++) {

projects/rccl/src/graph/topo.h

@@ -139,6 +139,12 @@ struct ncclTopoLinkList {
 #define RCCL_TOPO_FORCE_INTRA 16
 #define RCCL_TOPO_XGMI_ALL  32
 
+/* Rome preset graph: index into romeTopoModels[] in rome_models.cc, or sentinel below. */
+#define RCCL_ROME_TOPO_PRESET_MODEL_IDX_NONE (-1)
+#define RCCL_ROME_TOPO_PRESET_MODEL_IDX_GIO_COLUMBA (1000000)
+/* parse4H4P() applies rome_model_68 directly; this tags the preset, not romeTopoModels[]. */
+#define RCCL_ROME_TOPO_PRESET_MODEL_IDX_4H4P (1000001)
+
 
 #define GCN_ARCH_NAME_LEN 16
 
@@ -215,6 +221,8 @@ struct ncclTopoSystem {
   // [RCCL] Track hostIdx to support rail-optimized rings/trees
   int hostIdx;
   bool useRailOptimizedTrees;
+  /* RCCL Rome / GIO preset: RCCL_ROME_TOPO_PRESET_MODEL_IDX_* sentinels or romeTopoModels[] index */
+  int romeTopoModelIdx;
 };
 
 ncclResult_t ncclTopoGetNode(struct ncclTopoSystem* system, struct ncclTopoNode** node, int type, uint64_t id);

projects/rccl/src/init.cc

@@ -40,12 +40,14 @@
 #include <sys/resource.h>
 #include <unistd.h>
 #include "graph/topo.h"
+#include "graph/rome_topo_consensus.h"
 #include "graph/xml.h"
 #include "archinfo.h"
 #include "param.h"
 #include "nvtx_payload_schemas.h"
 #include "utils.h"
 #include <mutex>
+#include <unordered_map>
 #include "ce_coll.h"
 #include "nvtx.h"
 
@@ -1245,6 +1247,28 @@ static ncclResult_t initNvlDomainInfo(struct ncclComm* comm) {
   return ncclSuccess;
 }
 
+// True when every host has the same number of ranks and those partitions sum to nranks.
+static bool uniformRanksPerHost(const struct ncclComm* comm, int nranks) {
+  int ranksPerHost = -1;
+  int total = 0;
+  for (int i = 0; i < nranks; i++) {
+    uint64_t h = comm->peerInfo[i].hostHash;
+    bool lowestRankOnHost = true;
+    for (int j = 0; j < i; j++) {
+      if (comm->peerInfo[j].hostHash == h) { lowestRankOnHost = false; break; }
+    }
+    if (!lowestRankOnHost) continue;
+    int cnt = 0;
+    for (int j = 0; j < nranks; j++) {
+      if (comm->peerInfo[j].hostHash == h) cnt++;
+    }
+    total += cnt;
+    if (ranksPerHost < 0) ranksPerHost = cnt;
+    else if (cnt != ranksPerHost) return false;
+  }
+  return total == nranks && ranksPerHost > 0;
+}
+
 static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* parent, uint64_t timers[TIMERS_INIT_COUNT]) {
   // We use 2 AllGathers
   // 1. { peerInfo, comm, compCap}
@@ -1279,8 +1303,10 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
     int cpuVendor;
     int localRanks;
     int nc;
+    int romeTopoModelIdx;
     bool pivotA2AEnabled;
     bool ll128Enabled;
+    char hostname[128];
   };
 
   int nChannelsOrig;
@@ -1656,6 +1682,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
 
   allGather3Data[rank].cpuArch = comm->cpuArch;
   allGather3Data[rank].cpuVendor = comm->cpuVendor;
+  allGather3Data[rank].romeTopoModelIdx = comm->topo->romeTopoModelIdx;
+  (void) getHostName(allGather3Data[rank].hostname, sizeof(allGather3Data[rank].hostname), '\0');
 
   comm->nChannels = std::min(treeGraph->nChannels, ringGraph->nChannels);
 
@@ -1667,6 +1695,15 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
 
   NCCLCHECKGOTO(bootstrapAllGather(comm->bootstrap, allGather3Data, sizeof(*allGather3Data)), ret, fail);
 
+  if (uniformRanksPerHost(comm, nranks)) {
+    NCCLCHECKGOTO(rcclCheckRomeTopoModelIdxConsensus(
+        nranks,
+        [&](int r) { return allGather3Data[r].romeTopoModelIdx; },
+        [&](int r) { return allGather3Data[r].hostname; },
+        [&](int r) { return comm->peerInfo[r].hostHash; }),
+      ret, fail);
+  }
+
   // Determine nNodes, firstRanks, ...
   NCCLCHECKGOTO(ncclCalloc(&nodesFirstRank, nranks), ret, fail);
   NCCLCHECKGOTO(ncclCalloc(&nodesTreePatterns, nranks), ret, fail);

projects/rccl/test/CMakeLists.txt

@@ -199,6 +199,7 @@ if(BUILD_TESTS)
 
     set(TEST_FIXTURE_SOURCE_FILES
       BitOpsTests.cpp
+      RomeTopoConsensusTests.cpp
       CommTests.cpp
       EnqueueCountTests.cpp
       device/TestOp128.cpp

projects/rccl/test/RomeTopoConsensusTests.cpp

@@ -0,0 +1,62 @@
+/*************************************************************************
+ * Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#include "graph/rome_topo_consensus.h"
+#include "gtest/gtest.h"
+#include <array>
+
+namespace RcclUnitTesting {
+
+TEST(RomeTopoConsensus, emptyRanks) {
+  EXPECT_EQ(rcclCheckRomeTopoModelIdxConsensus(
+      0,
+      [](int) { return 0; },
+      [](int) { return ""; },
+      [](int) { return 0ULL; }),
+    ncclSuccess);
+}
+
+TEST(RomeTopoConsensus, allSameModelIdx) {
+  constexpr int n = 4;
+  std::array<int, n> idx{{7, 7, 7, 7}};
+  std::array<const char*, n> names{{"a", "a", "b", "b"}};
+  std::array<uint64_t, n> hosts{{1, 1, 2, 2}};
+  EXPECT_EQ(rcclCheckRomeTopoModelIdxConsensus(
+      n,
+      [&](int r) { return idx[r]; },
+      [&](int r) { return names[r]; },
+      [&](int r) { return hosts[r]; }),
+    ncclSuccess);
+}
+
+TEST(RomeTopoConsensus, mismatchFails) {
+  constexpr int n = 2;
+  std::array<int, n> idx{{1, 2}};
+  std::array<const char*, n> names{{"h0", "h1"}};
+  std::array<uint64_t, n> hosts{{10, 11}};
+  EXPECT_EQ(rcclCheckRomeTopoModelIdxConsensus(
+      n,
+      [&](int r) { return idx[r]; },
+      [&](int r) { return names[r]; },
+      [&](int r) { return hosts[r]; }),
+    ncclInvalidUsage);
+}
+
+TEST(RomeTopoConsensus, tieBreakLowerFirstRankWinsThenMinorityFails) {
+  // idx 2 at ranks 0,1; idx 3 at ranks 2,3 -> plurality tie; first occurrence rank 0 -> refIdx 2
+  constexpr int n = 4;
+  std::array<int, n> idx{{2, 2, 3, 3}};
+  std::array<const char*, n> names{{"a", "a", "a", "a"}};
+  std::array<uint64_t, n> hosts{{1, 1, 1, 1}};
+  EXPECT_EQ(rcclCheckRomeTopoModelIdxConsensus(
+      n,
+      [&](int r) { return idx[r]; },
+      [&](int r) { return names[r]; },
+      [&](int r) { return hosts[r]; }),
+    ncclInvalidUsage);
+}
+
+} // namespace RcclUnitTesting

projects/rccl/tools/topo_expl/include/topo_expl_impl.h

@@ -25,6 +25,7 @@ struct allGatherInfo {
   struct graphInfo graphInfo[NCCL_NUM_ALGORITHMS];
   struct ncclTopoRanks topoRanks;
   int nc;
+  int romeTopoModelIdx;
   bool pivotA2AEnabled;
   bool ll128Enabled;
   bool mscclEnabled;