fix(rccl): index topology GPUs by SMI index, not HIP ordinal

projects/rccl/CHANGELOG.md

@@ -48,6 +48,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it.
 
 ### Resolved issues
+* Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, causing an `ncclGather` performance regression on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU.
 * Fixed profiler overhead when no profiler plugin is loaded or only kernel-channel events are enabled.
 * Fixed profiler API events reporting rank 0 instead of the originating communicator rank.
 * Fixed P2P IPC registration reuse producing out-of-bounds remote addresses when a registered allocation spans multiple cuMem segments.

projects/rccl/src/graph/paths.cc

@@ -1191,12 +1191,12 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm*
 
   // detect if all GPUs are connected by XGMI
   for (int i = 0; i < system->nodes[GPU].count && allXgmi; i++) {
-    int cudaDev1 = system->nodes[GPU].nodes[i].gpu.dev;
+    int dev1 = system->nodes[GPU].nodes[i].gpu.dev;
     for (int j = 0; j < system->nodes[GPU].count && allXgmi; j++) {
       if (i == j) continue;
-      int cudaDev2 = system->nodes[GPU].nodes[j].gpu.dev;
+      int dev2 = system->nodes[GPU].nodes[j].gpu.dev;
       bool isXGMI;
-      NCCLCHECKGOTO(ncclTopoGetLinkType(comm->topo, cudaDev1, cudaDev2, &isXGMI), ret, fail);
+      NCCLCHECKGOTO(ncclTopoGetLinkType(comm->topo, dev1, dev2, &isXGMI), ret, fail);
       allXgmi &= isXGMI;
     }
   }

projects/rccl/src/graph/search.cc

@@ -1743,17 +1743,17 @@ ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, str
   return ncclSuccess;
 }
 
-ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI, int maxInter,
+ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int dev1, int dev2, bool* isXGMI, int maxInter,
                                  int nInter, int* inter) {
   int interGpus[MAX_XGMI_INTER_GPUS + 1];
   int ngpus = system->nodes[GPU].count;
   *isXGMI = false;
   // check for direct XGMI connection
   for (int i = 0; i < ngpus; i++) {
-    if (system->nodes[GPU].nodes[i].gpu.dev == cudaDev1) {
+    if (system->nodes[GPU].nodes[i].gpu.dev == dev1) {
       struct ncclTopoNode* node = system->nodes[GPU].nodes + i;
       for (int k = 0; k < system->nodes[GPU].count; k++) {
-        if (system->nodes[GPU].nodes[k].gpu.dev == cudaDev2) {
+        if (system->nodes[GPU].nodes[k].gpu.dev == dev2) {
           // NCCL 2.30 routes GPU->GPU through DEV nodes, so detect XGMI via path type (not a single hop).
           if (node->paths[GPU][k].count > 0 && node->paths[GPU][k].type == PATH_NVL) {
             *isXGMI = true;
@@ -1775,27 +1775,27 @@ ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, in
     }
     if (j < nInter) return ncclSuccess;
     if (nInter > 0 && inter != nullptr) {
-      ncclTopoGetLinkType(system, inter[nInter], cudaDev2, &res2, 0);
+      ncclTopoGetLinkType(system, inter[nInter], dev2, &res2, 0);
       if (res2) {
         *isXGMI = true;
         return ncclSuccess;
       }
       memcpy(interGpus + 1, inter + 1, sizeof(int) * nInter);
     }
-    interGpus[0] = cudaDev1;
+    interGpus[0] = dev1;
     // add one more intermediate GPU recursively util reaching max depth
     nInter++;
     if (nInter + 2 > ngpus || nInter > MAX_XGMI_INTER_GPUS || nInter > maxInter) return ncclSuccess;
     for (int i = 0; i < ngpus; i++) {
       int dev = system->nodes[GPU].nodes[i].gpu.dev;
       // skip duplicated GPU
-      if (dev == cudaDev2) continue;
+      if (dev == dev2) continue;
       for (j = 0; j < nInter; j++)
         if (dev == interGpus[j]) break;
       if (j < nInter) continue;
       // check connectivity with intermediate GPUs
       interGpus[nInter] = dev;
-      ncclTopoGetLinkType(system, cudaDev1, cudaDev2, &res3, maxInter, nInter, interGpus);
+      ncclTopoGetLinkType(system, dev1, dev2, &res3, maxInter, nInter, interGpus);
       if (res3) {
         *isXGMI = true;
         return ncclSuccess;

projects/rccl/src/graph/xml.cc

@@ -952,7 +952,7 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml*
 // at that BDF is missing or a different function. Link those HIP siblings with
 // 1-hop XGMI even when SMI only enumerates the physical GPU.
 static ncclResult_t ncclTopoXmlAddHipPartitionXgmi(struct ncclXml* xml, struct ncclXmlNode* gpuNode,
-                                                   struct ncclXmlNode* pciNode, uint32_t xmlDev) {
+                                                   struct ncclXmlNode* pciNode, uint32_t hipDev) {
   const char* pciBusId = NULL;
   NCCLCHECK(xmlGetAttr(pciNode, "busid", &pciBusId));
   if (pciBusId == NULL) return ncclSuccess;
@@ -962,7 +962,7 @@ static ncclResult_t ncclTopoXmlAddHipPartitionXgmi(struct ncclXml* xml, struct n
   int nHip = 0;
   if (hipGetDeviceCount(&nHip) != hipSuccess) return ncclSuccess;
   for (int h = 0; h < nHip; h++) {
-    if (h == (int)xmlDev) continue;
+    if (h == (int)hipDev) continue;
     char peerBus[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];
     if (hipDeviceGetPCIBusId(peerBus, sizeof(peerBus), h) != hipSuccess) continue;
     int64_t peerId = 0;
@@ -985,18 +985,20 @@ static ncclResult_t ncclTopoXmlAddHipPartitionXgmi(struct ncclXml* xml, struct n
 }
 #endif
 
-ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t xmlDev, uint32_t smiDev, struct ncclXml* xml,
+ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t hipDev, uint32_t smiDev, struct ncclXml* xml,
                                    struct ncclXmlNode** gpuNodeRet) {
   struct ncclXmlNode* gpuNode = NULL;
-  NCCLCHECK(xmlGetSubKvInt(pciNode, "gpu", &gpuNode, "dev", (int)xmlDev));
+  // Processes learn about each other's GPUs only through XML fusion, so "dev" must mean the same in all
+  // of them. hipDev depends on each process's device visibility; the SMI index does not.
+  NCCLCHECK(xmlGetSubKvInt(pciNode, "gpu", &gpuNode, "dev", (int)smiDev));
   if (gpuNode == NULL) NCCLCHECK(xmlAddNode(xml, pciNode, "gpu", &gpuNode));
 
   int index = -1;
 
   int dev = -1;
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "dev", &index));
   if (index == -1) {
-    NCCLCHECK(xmlSetAttrInt(gpuNode, "dev", xmlDev));
+    NCCLCHECK(xmlSetAttrInt(gpuNode, "dev", smiDev));
   }
   NCCLCHECK(xmlGetAttrInt(gpuNode, "dev", &dev));
   if (dev == -1) {
@@ -1007,7 +1009,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t xmlDev,
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "sm", &index));
   if (index == -1) {
     cudaDeviceProp devProp;
-    CUDACHECK(cudaGetDeviceProperties(&devProp, (int)xmlDev));
+    CUDACHECK(cudaGetDeviceProperties(&devProp, (int)hipDev));
     NCCLCHECK(xmlSetAttrInt(gpuNode, "sm", devProp.multiProcessorCount));
   }
   int sm;
@@ -1018,7 +1020,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t xmlDev,
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "gcn", &index));
   if (index == -1) {
     hipDeviceProp_t devProp;
-    CUDACHECK(hipGetDeviceProperties(&devProp, (int)xmlDev));
+    CUDACHECK(hipGetDeviceProperties(&devProp, (int)hipDev));
     // extract only the releveant info from the gcnArchName attribute
     // e.g.: convert "gfx908:sramecc+:xnack-" to "gfx908"
     char gcnArchNameSubstr[128];
@@ -1036,7 +1038,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t xmlDev,
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "arch", &index));
   if (index == -1) {
     hipDeviceProp_t devProp;
-    CUDACHECK(hipGetDeviceProperties(&devProp, (int)xmlDev));
+    CUDACHECK(hipGetDeviceProperties(&devProp, (int)hipDev));
     memcpy(&arch.arch, &devProp.arch, sizeof(hipDeviceArch_t));
     NCCLCHECK(xmlSetAttrInt(gpuNode, "arch", arch.value));
   }
@@ -1240,7 +1242,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t xmlDev,
 #endif
   }
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
-  NCCLCHECK(ncclTopoXmlAddHipPartitionXgmi(xml, gpuNode, pciNode, xmlDev));
+  NCCLCHECK(ncclTopoXmlAddHipPartitionXgmi(xml, gpuNode, pciNode, hipDev));
 #endif
 #if CUDART_VERSION >= 11080
   struct ncclXmlNode* c2cNode = NULL;
@@ -1373,7 +1375,7 @@ ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct nccl
   NCCLCHECK(rocm_smi_getDeviceIndexByPciBusId(busId, &smiIndex));
 #endif
   int hipDev = -1;
-  if (hipDeviceGetByPCIBusId(&hipDev, busId) != hipSuccess || hipDev < 0) hipDev = (int)smiIndex;
+  CUDACHECK(hipDeviceGetByPCIBusId(&hipDev, busId));
   NCCLCHECK(ncclTopoGetXmlFromGpu(node, (uint32_t)hipDev, smiIndex, xml, gpuNode));
 #else
   nvmlDevice_t nvmlDev;

projects/rccl/src/include/graph.h

@@ -86,7 +86,7 @@ ncclResult_t ncclGetUserP2pLevel(int* level);
 #define MAX_XGMI_INTER_GPUS 4
 ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, struct ncclTopoGraph* graph, int channelId,
                                     int type, int64_t* id, int* dev);
-ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI,
+ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int dev1, int dev2, bool* isXGMI,
                                  int maxInter = MAX_XGMI_INTER_GPUS, int nInter = 0, int* inter = nullptr);
 
 // Find CPU affinity

projects/rccl/src/init.cc

@@ -2148,7 +2148,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
 
       bool isXGMI;
       // Limit to single intermediate GPU for enabling clique
-      NCCLCHECK(ncclTopoGetLinkType(comm->topo, i, j, &isXGMI, 1));
+      NCCLCHECK(ncclTopoGetLinkType(comm->topo, comm->peerInfo[i].nvmlDev, comm->peerInfo[j].nvmlDev, &isXGMI, 1));
       allXgmi &= isXGMI;
     }
   }

projects/rccl/src/transport.cc

@@ -34,7 +34,7 @@ static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph*
     NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, peer, graph, channelId, (type == 1) ? 0 : 1, nullptr, &n2));
   }
   bool xgmi;
-  NCCLCHECK(ncclTopoGetLinkType(comm->topo, myInfo->cudaDev, peerInfo->cudaDev, &xgmi));
+  NCCLCHECK(ncclTopoGetLinkType(comm->topo, myInfo->nvmlDev, peerInfo->nvmlDev, &xgmi));
 
   for (int t = 0; t < NTRANSPORTS; t++) {
     if (graph == NULL && connIndex == NCCL_CONN_IDX_P2P_NET && (t == TRANSPORT_SHM || (!xgmi && t == TRANSPORT_P2P)))

projects/rccl/src/transport/p2p.cc

@@ -471,7 +471,7 @@ ncclResult_t p2pSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st
 
   resources->next_hdp_reg = 0;
   bool isXGMI;
-  if (ncclTopoGetLinkType(comm->topo, myInfo->cudaDev, peerInfo->cudaDev, &isXGMI) != ncclSuccess) {
+  if (ncclTopoGetLinkType(comm->topo, myInfo->nvmlDev, peerInfo->nvmlDev, &isXGMI) != ncclSuccess) {
     INFO(NCCL_INIT | NCCL_P2P, "Ring %02d : %d -> %d failed to get link type and hop count", channelId, myInfo->rank,
          peerInfo->rank);
     return ncclInternalError;