From fa7a9488704f4f71a18f20c0747cda44763fb02f Mon Sep 17 00:00:00 2001
From: Anatolii Rozanov <arozanov@amd.com>
Date: Thu, 15 Oct 2026 12:32:11 +0200
Subject: [PATCH 1/7] fix(rccl): index topology GPUs by SMI index, not HIP
 ordinal

Since #10640 the topology XML stamps each GPU's "dev" with the
process-local HIP ordinal. On MI355X the HIP order differs from the
PCI/SMI order, so model matching, ring and tree order and NIC
assignment changed, raising multi-node ncclGather latency for 16 KB
to 1 MB by 10-17%. When each rank sees only its own GPU, every GPU
gets dev 0 and communicator initialization fails.

Stamp "dev" with the SMI index again and keep the HIP ordinal for the
HIP property queries. ncclTopoGetLinkType matches its arguments
against gpu.dev, so the transport and P2P callers now pass nvmlDev,
and its parameters are renamed from cudaDev1/cudaDev2 to dev1/dev2.

Co-authored-by: Cursor <cursoragent@cursor.com>
---
 projects/rccl/CHANGELOG.md                   |  1 +
 projects/rccl/src/graph/paths.cc             |  6 +--
 projects/rccl/src/graph/search.cc            | 14 +++---
 projects/rccl/src/graph/xml.cc               | 20 ++++----
 projects/rccl/src/include/graph.h            |  2 +-
 projects/rccl/src/transport.cc               |  2 +-
 projects/rccl/src/transport/p2p.cc           |  2 +-
 projects/rccl/test/CMakeLists.txt            |  1 +
 projects/rccl/test/graph/TopoGpuDevTests.cpp | 50 ++++++++++++++++++++
 9 files changed, 76 insertions(+), 22 deletions(-)
 create mode 100644 projects/rccl/test/graph/TopoGpuDevTests.cpp

diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md
index db230746d30a..f5aa3a994cd1 100644
--- a/projects/rccl/CHANGELOG.md
+++ b/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, raising 16 KB to 1 MB `ncclGather` latency by up to 18% 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.
diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc
index 66e465d2cbe8..fcbb689f0c6f 100644
--- a/projects/rccl/src/graph/paths.cc
+++ b/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;
     }
   }
diff --git a/projects/rccl/src/graph/search.cc b/projects/rccl/src/graph/search.cc
index 52920b56314a..867ff9938678 100644
--- a/projects/rccl/src/graph/search.cc
+++ b/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;
diff --git a/projects/rccl/src/graph/xml.cc b/projects/rccl/src/graph/xml.cc
index dfb98409da38..dd09a4a65e36 100644
--- a/projects/rccl/src/graph/xml.cc
+++ b/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,10 +985,12 @@ 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;
@@ -996,7 +998,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t xmlDev,
   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;
diff --git a/projects/rccl/src/include/graph.h b/projects/rccl/src/include/graph.h
index 4c7484d67c6c..a65b10002714 100644
--- a/projects/rccl/src/include/graph.h
+++ b/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
diff --git a/projects/rccl/src/transport.cc b/projects/rccl/src/transport.cc
index 8d6f69cdd6e0..2f47b046c727 100644
--- a/projects/rccl/src/transport.cc
+++ b/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)))
diff --git a/projects/rccl/src/transport/p2p.cc b/projects/rccl/src/transport/p2p.cc
index 0b0c6391f6b6..e88e131aa2a1 100644
--- a/projects/rccl/src/transport/p2p.cc
+++ b/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;
diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt
index 84e40b6dd88a..b708a738cb05 100644
--- a/projects/rccl/test/CMakeLists.txt
+++ b/projects/rccl/test/CMakeLists.txt
@@ -790,6 +790,7 @@ if(BUILD_TESTS)
       graph/XmlTests.cpp
       graph/NetDevsPolicyP2pNetTests.cpp
       graph/P2pPxnLevel1NetIdTests.cpp
+      graph/TopoGpuDevTests.cpp
       graph/TopoTests.cpp
       transport/NetIbCast/NetIbCastPlaneSubnetTests.cpp
       transport/NetIbCast/NetIbCastMultiplaneTests.cpp
diff --git a/projects/rccl/test/graph/TopoGpuDevTests.cpp b/projects/rccl/test/graph/TopoGpuDevTests.cpp
new file mode 100644
index 000000000000..43d81dba4b7c
--- /dev/null
+++ b/projects/rccl/test/graph/TopoGpuDevTests.cpp
@@ -0,0 +1,50 @@
+/*************************************************************************
+ * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+// Checks that the GPU nodes of a real topology carry the SMI index as `dev`, distinct per GPU.
+// Needs at least two GPUs; on a host whose HIP order matches the SMI order it cannot fail.
+
+#include <gtest/gtest.h>
+#include <rccl/rccl.h>
+
+#include <set>
+#include <vector>
+
+#include "comm.h"
+#include "graph/topo.h"
+#include "../common/ProcessIsolatedTestRunner.hpp"
+
+namespace RcclUnitTesting {
+
+TEST(TopoGpuDev, OwnGpuDevIsNvmlDevAndUnique)
+{
+  RUN_ISOLATED_TEST("TopoGpuDev_OwnGpuDevIsNvmlDevAndUnique", []() {
+    int numDevices = 0;
+    ASSERT_EQ(hipGetDeviceCount(&numDevices), hipSuccess);
+    if (numDevices < 2) GTEST_SKIP() << "Needs at least 2 GPUs";
+
+    std::vector<ncclComm_t> comms(numDevices, nullptr);
+    ASSERT_EQ(ncclCommInitAll(comms.data(), numDevices, nullptr), ncclSuccess);
+
+    for (ncclComm_t comm : comms) {
+      const struct ncclTopoNodeSet& gpus = comm->topo->nodes[GPU];
+      int ownDev = -1;
+      for (int gpuIndex = 0; gpuIndex < gpus.count; gpuIndex++) {
+        if (gpus.nodes[gpuIndex].gpu.rank == comm->rank) ownDev = gpus.nodes[gpuIndex].gpu.dev;
+      }
+      EXPECT_EQ(ownDev, comm->nvmlDev) << "rank " << comm->rank << " (HIP device " << comm->cudaDev << ")";
+    }
+
+    const struct ncclTopoNodeSet& gpus = comms[0]->topo->nodes[GPU];
+    std::set<int> devs;
+    for (int gpuIndex = 0; gpuIndex < gpus.count; gpuIndex++) devs.insert(gpus.nodes[gpuIndex].gpu.dev);
+    EXPECT_EQ(static_cast<int>(devs.size()), gpus.count) << "GPU nodes share a dev index";
+
+    for (ncclComm_t comm : comms) EXPECT_EQ(ncclCommDestroy(comm), ncclSuccess);
+  });
+}
+
+} // namespace RcclUnitTesting

From 48206137a928a2eec9d5a86d6d4bd98b793a883c Mon Sep 17 00:00:00 2001
From: Anatolii Rozanov <arozanov@amd.com>
Date: Thu, 15 Oct 2026 13:56:17 +0200
Subject: [PATCH 2/7] fix(rccl): look up the XGMI clique by SMI index

The clique check in initTransportsRank passed rank indices to
ncclTopoGetLinkType, which matches its arguments against gpu.dev, the
SMI index. With a GPU subset such as HIP_VISIBLE_DEVICES=2,3 the
lookup misses both GPUs and allXgmi turns false, which halves the
gfx906 channel multiplier. Pass each peer's nvmlDev instead.

Co-authored-by: Cursor <cursoragent@cursor.com>
---
 projects/rccl/src/init.cc | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc
index 2f4da2afc47f..cd009eb7623c 100644
--- a/projects/rccl/src/init.cc
+++ b/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;
     }
   }

From acbe1c7c6f227faaaa5297a3a14e4913fe6ccb9c Mon Sep 17 00:00:00 2001
From: Anatolii Rozanov <arozanov@amd.com>
Date: Thu, 15 Oct 2026 14:59:10 +0200
Subject: [PATCH 3/7] test(rccl): run TopoGpuDev in the fixtures-debug test
 lists

rccl-UnitTestsFixturesDebug is run with explicit filters, so the new
TopoGpuDev.* tests only run where they are listed. Add them next to
TopoTest.* in the CTest categories and in the test_runner plans that
already run TopoTest.*.

Co-authored-by: Cursor <cursoragent@cursor.com>
---
 projects/rccl/test/test_categories_fixtures_debug.yaml       | 1 +
 .../scripts/test_runner/configs/mi300x_mellanox_ib.json      | 5 +++++
 .../tools/scripts/test_runner/configs/mi355x_thor2_roce.json | 5 +++++
 .../tools/scripts/test_runner/configs/mi455_ainic_roce.json  | 5 +++++
 .../scripts/test_runner/configs/mi455_ainic_roce_func.json   | 5 +++++
 5 files changed, 21 insertions(+)

diff --git a/projects/rccl/test/test_categories_fixtures_debug.yaml b/projects/rccl/test/test_categories_fixtures_debug.yaml
index 5769653bdaef..0e9874c498ad 100644
--- a/projects/rccl/test/test_categories_fixtures_debug.yaml
+++ b/projects/rccl/test/test_categories_fixtures_debug.yaml
@@ -54,6 +54,7 @@ test_categories:
       - "P2pPxnLevel1NetId.*"
       - "SymmetricMemP2pLevelTests.*"
       - "TopoTest.*"
+      - "TopoGpuDev.*"
       - "TopoNicFusionTests.*"
       - "P2pMaxNchannels.*"
       - "ChannelDefaults.*"
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json
index fda03570f133..47fe2814f84c 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json
@@ -929,6 +929,11 @@
           "description": "Topology construction tests",
           "test_filter": "TopoTest.*"
         },
+        {
+          "name": "TopoGpuDev",
+          "description": "Topology GPU dev index tests",
+          "test_filter": "TopoGpuDev.*"
+        },
         {
           "name": "P2pPxnLevel1NetId",
           "description": "PXN level-1 packed netId lookup (NCCL PR #2258 / AICOMRCCL-2038)",
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json
index ef63799e08d3..93b6458d08ed 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json
@@ -696,6 +696,11 @@
           "description": "Topology construction tests",
           "test_filter": "TopoTest.*"
         },
+        {
+          "name": "TopoGpuDev",
+          "description": "Topology GPU dev index tests",
+          "test_filter": "TopoGpuDev.*"
+        },
         {
           "name": "P2pPxnLevel1NetId",
           "description": "PXN level-1 packed netId lookup (NCCL PR #2258 / AICOMRCCL-2038)",
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json
index 0e315dd71608..83959e389a6f 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json
@@ -1424,6 +1424,11 @@
           "description": "Topology construction tests",
           "test_filter": "TopoTest.*"
         },
+        {
+          "name": "TopoGpuDev",
+          "description": "Topology GPU dev index tests",
+          "test_filter": "TopoGpuDev.*"
+        },
         {
           "name": "P2pPxnLevel1NetId",
           "description": "PXN level-1 packed netId lookup (NCCL PR #2258 / AICOMRCCL-2038)",
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json
index 26748e8c8ba0..e24d456638e8 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json
@@ -1483,6 +1483,11 @@
           "description": "Topology construction tests",
           "test_filter": "TopoTest.*"
         },
+        {
+          "name": "TopoGpuDev",
+          "description": "Topology GPU dev index tests",
+          "test_filter": "TopoGpuDev.*"
+        },
         {
           "name": "P2pPxnLevel1NetId",
           "description": "PXN level-1 packed netId lookup (NCCL PR #2258 / AICOMRCCL-2038)",

From a504859db7f72c85e91dead2a917b79635595f13 Mon Sep 17 00:00:00 2001
From: Anatolii Rozanov <arozanov@amd.com>
Date: Thu, 15 Oct 2026 16:25:40 +0200
Subject: [PATCH 4/7] fix(rccl): fail topology fill if the GPU has no HIP
 ordinal

ncclTopoFillGpu fell back to the SMI index when hipDeviceGetByPCIBusId
failed and then passed it to hipGetDeviceProperties as a HIP ordinal,
which queries an unrelated device or fails later with a misleading
error. The GPU being filled is always the caller's own, so the lookup
cannot legitimately fail: report the HIP error instead.

Also correct the CHANGELOG figure to the 9-15% measured in the 6-round
runs.

Co-authored-by: Cursor <cursoragent@cursor.com>
---
 projects/rccl/CHANGELOG.md     | 2 +-
 projects/rccl/src/graph/xml.cc | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md
index f5aa3a994cd1..365032cb638c 100644
--- a/projects/rccl/CHANGELOG.md
+++ b/projects/rccl/CHANGELOG.md
@@ -48,7 +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, raising 16 KB to 1 MB `ncclGather` latency by up to 18% on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU.
+* Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, raising 16 KB to 1 MB `ncclGather` latency by up to 15% 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.
diff --git a/projects/rccl/src/graph/xml.cc b/projects/rccl/src/graph/xml.cc
index dd09a4a65e36..561ab85ab74d 100644
--- a/projects/rccl/src/graph/xml.cc
+++ b/projects/rccl/src/graph/xml.cc
@@ -1375,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;

From 067c488450d161bffc2bfcd8c129b28d6b287df4 Mon Sep 17 00:00:00 2001
From: Anatolii Rozanov <arozanov@amd.com>
Date: Thu, 15 Oct 2026 16:25:40 +0200
Subject: [PATCH 5/7] test(rccl): make the GPU dev tests fail on any multi-GPU
 host

TopoGpuDev only failed where HIP order already differs from SMI order,
so it could not catch the regression on CI. Run it a second time with
HIP_VISIBLE_DEVICES=1,0: on any host with two or more GPUs one of the
two runs maps HIP ordinals to different SMI indices.

Neither test covered per-process visibility, where every process sees
its own GPU as ordinal 0. Generalize AsymmetricVisibility to a layout
of N workers and add OneGpuPerRankAllReduce, one rank per GPU with
only that GPU visible. Two ranks are not enough to break develop;
eight are. List the new test in the test_runner plans that already
run AsymmetricVisibility.

Co-authored-by: Cursor <cursoragent@cursor.com>
---
 .../rccl/test/AsymmetricVisibilityTests.cpp   | 124 +++++++++++-------
 projects/rccl/test/graph/TopoGpuDevTests.cpp  |  62 +++++----
 .../configs/mi300x_mellanox_ib.json           |   3 +-
 .../configs/mi355x_thor2_roce.json            |   5 +
 .../test_runner/configs/mi455_ainic_roce.json |   5 +
 .../configs/mi455_ainic_roce_func.json        |   5 +
 6 files changed, 129 insertions(+), 75 deletions(-)

diff --git a/projects/rccl/test/AsymmetricVisibilityTests.cpp b/projects/rccl/test/AsymmetricVisibilityTests.cpp
index 213fbde67bd2..b2891d6b917a 100644
--- a/projects/rccl/test/AsymmetricVisibilityTests.cpp
+++ b/projects/rccl/test/AsymmetricVisibilityTests.cpp
@@ -17,8 +17,8 @@
 // This test spawns its own worker processes via fork()+execv(), each with a
 // distinct HIP_VISIBLE_DEVICES set before HIP is initialized.
 // The orchestrator generates the ncclUniqueId and hands it to the workers
-// hex-encoded in the environment. Each worker binds its device, joins a
-// 2-rank communicator, asserts that init left no pending HIP error, and runs a
+// hex-encoded in the environment. Each worker binds its device, joins the
+// communicator, asserts that init left no pending HIP error, and runs a
 // small AllReduce to confirm the communicator is functional.
 //
 // Note: the legacy IPC path requires the cuMem/VMM handle path to survive
@@ -38,21 +38,41 @@
 #include <cstdlib>
 #include <cstring>
 #include <string>
+#include <vector>
 
 namespace RcclUnitTesting {
 namespace {
 
 // Env var names shared between orchestrator and worker.
-constexpr const char* kEnvRank = "RCCL_ASYM_RANK"; // worker rank (0 or 1)
+constexpr const char* kEnvRank = "RCCL_ASYM_RANK"; // worker rank
+constexpr const char* kEnvNumRanks = "RCCL_ASYM_NRANKS";
 constexpr const char* kEnvLocalDev =
   "RCCL_ASYM_LOCALDEV"; // ordinal to bind within visible set
 constexpr const char* kEnvUid = "RCCL_ASYM_UID"; // hex-encoded ncclUniqueId
 
-constexpr int kNumRanks = 2;
+// Which devices a worker sees and which ordinal within them it binds.
+struct WorkerSpec {
+  std::string visibleDevices;
+  int localDev;
+};
+
+// A layout returns one WorkerSpec per rank, or nothing if numDevices is too few.
+using VisibilityLayout = std::vector<WorkerSpec> (*)(int numDevices);
 
-// Minimum GPUs for the asymmetric topology: rank0 sees 2 devices, rank1 sees
-// a third, disjoint device.
-constexpr int kMinDevices = 3;
+// ROCM-27034: rank0 sees devices {0,1} and binds ordinal 1, rank1 sees {2}.
+std::vector<WorkerSpec> asymmetricLayout(int numDevices) {
+  if (numDevices < 3) { return {}; }
+  return {{"0,1", 1}, {"2", 0}};
+}
+
+// One rank per GPU, each seeing only its own GPU, so every rank binds HIP
+// ordinal 0 and the topology has to tell the GPUs apart by something else.
+std::vector<WorkerSpec> oneGpuPerRankLayout(int numDevices) {
+  if (numDevices < 2) { return {}; }
+  std::vector<WorkerSpec> workers;
+  for (int dev = 0; dev < numDevices; ++dev) { workers.push_back({std::to_string(dev), 0}); }
+  return workers;
+}
 
 // Process exit codes shared between workers and coordinator. kCoordSkip is
 // returned by the coordinator when there aren't enough GPUs; the parent maps
@@ -138,7 +158,7 @@ struct WorkerResources {
   } while (0)
 
 // Body executed inside a re-exec'd worker process. Returns kOk on success.
-int runWorker(int rank, int localDev, const ncclUniqueId& id) {
+int runWorker(int rank, int numRanks, int localDev, const ncclUniqueId& id) {
   char tag[32];
   snprintf(tag, sizeof(tag), "[asym rank%d]", rank);
 
@@ -151,7 +171,7 @@ int runWorker(int rank, int localDev, const ncclUniqueId& id) {
 
   WorkerResources res;
   ncclResult_t nres;
-  ASYM_NCCLCHECK(ncclCommInitRank(&res.comm, kNumRanks, id, rank), tag,
+  ASYM_NCCLCHECK(ncclCommInitRank(&res.comm, numRanks, id, rank), tag,
                  return kWorkerFailed);
 
   // Core of the ROCM-27034 regression: a successful ncclCommInitRank must not
@@ -186,7 +206,7 @@ int runWorker(int rank, int localDev, const ncclUniqueId& id) {
                           hipMemcpyDeviceToHost),
                 tag, return kWorkerFailed);
 
-  const float expected = static_cast<float>(kNumRanks) * (kNumRanks + 1) / 2.0f;
+  const float expected = static_cast<float>(numRanks) * (numRanks + 1) / 2.0f;
   if (result != expected) {
     fprintf(stderr, "%s wrong AllReduce result %.1f (expected %.1f)\n", tag,
             result, expected);
@@ -198,7 +218,7 @@ int runWorker(int rank, int localDev, const ncclUniqueId& id) {
 }
 
 // Spawn one worker: fork, set the asymmetric environment, re-exec the worker.
-pid_t spawnWorker(int rank, const char* visibleDevices, int localDev,
+pid_t spawnWorker(int rank, int numRanks, const WorkerSpec& spec,
                   const std::string& uidHex) {
   pid_t pid = fork();
   if (pid != 0) {
@@ -207,10 +227,11 @@ pid_t spawnWorker(int rank, const char* visibleDevices, int localDev,
 
   // Child: set the per-rank environment BEFORE HIP initializes, then re-exec
   // a fresh image so the restricted visibility takes effect cleanly.
-  setenv("HIP_VISIBLE_DEVICES", visibleDevices, 1);
-  setenv("CUDA_VISIBLE_DEVICES", visibleDevices, 1);
+  setenv("HIP_VISIBLE_DEVICES", spec.visibleDevices.c_str(), 1);
+  setenv("CUDA_VISIBLE_DEVICES", spec.visibleDevices.c_str(), 1);
   setenv(kEnvRank, std::to_string(rank).c_str(), 1);
-  setenv(kEnvLocalDev, std::to_string(localDev).c_str(), 1);
+  setenv(kEnvNumRanks, std::to_string(numRanks).c_str(), 1);
+  setenv(kEnvLocalDev, std::to_string(spec.localDev).c_str(), 1);
   setenv(kEnvUid, uidHex.c_str(), 1);
   setenv("NCCL_CUMEM_ENABLE", "1", 0);
   // The worker filter lands on shard 0 only; inherited shard vars make it run 0 tests and exit kOk.
@@ -233,43 +254,41 @@ pid_t spawnWorker(int rank, const char* visibleDevices, int localDev,
 // fresh child for every collective test, and HIP is not fork-safe once
 // initialized in the parent, so touching it here would crash later TestBed
 // children. Returns an exit code the parent maps to skip/pass/fail.
-int runCoordinator() {
+int runCoordinator(VisibilityLayout layout) {
   const char* tag = "[asym]";
 
   int numDevices = 0;
   hipError_t err;
   ASYM_HIPCHECK(hipGetDeviceCount(&numDevices), tag, return kWorkerFailed);
-  if (numDevices < kMinDevices) { return kCoordSkip; }
+  const std::vector<WorkerSpec> workers = layout(numDevices);
+  if (workers.empty()) { return kCoordSkip; }
+  const int numRanks = static_cast<int>(workers.size());
 
   ncclUniqueId id;
   ncclResult_t nres;
   ASYM_NCCLCHECK(ncclGetUniqueId(&id), tag, return kWorkerFailed);
   const std::string uidHex = toHex(id);
 
-  // rank0: HIP_VISIBLE_DEVICES=0,1 -> bind ordinal 1; rank1: =2 -> bind ordinal 0.
-  const pid_t child0 = spawnWorker(0, "0,1", 1, uidHex);
-  const pid_t child1 = (child0 > 0) ? spawnWorker(1, "2", 0, uidHex) : -1;
-  if (child0 <= 0 || child1 <= 0) {
-    if (child0 > 0) {
-      kill(child0, SIGKILL);
-      waitpid(child0, nullptr, 0);
+  std::vector<pid_t> children;
+  for (int rank = 0; rank < numRanks; ++rank) {
+    const pid_t child = spawnWorker(rank, numRanks, workers[rank], uidHex);
+    if (child <= 0) {
+      fprintf(stderr, "%s fork of rank %d failed\n", tag, rank);
+      for (pid_t spawned : children) {
+        kill(spawned, SIGKILL);
+        waitpid(spawned, nullptr, 0);
+      }
+      return kWorkerFailed;
     }
-    if (child1 > 0) {
-      kill(child1, SIGKILL);
-      waitpid(child1, nullptr, 0);
-    }
-    fprintf(stderr, "%s fork failed (child0=%d, child1=%d)\n", tag, child0,
-            child1);
-    return kWorkerFailed;
+    children.push_back(child);
   }
 
-  // Reap whichever worker exits first. If it failed, kill its peer right
-  // away instead of waiting on it - a rank already blocked in ncclAllReduce
+  // Reap workers as they exit. Once one fails, kill the rest right away
+  // instead of waiting on them - a rank already blocked in ncclAllReduce
   // would otherwise hang forever for a partner that already gave up.
-  const pid_t children[kNumRanks] = {child0, child1};
-  bool reaped[kNumRanks] = {false, false};
+  std::vector<bool> reaped(numRanks, false);
   int rc = kOk;
-  for (int done_count = 0; done_count < kNumRanks; ++done_count) {
+  for (int done_count = 0; done_count < numRanks; ++done_count) {
     int status = 0;
     const pid_t done = waitpid(-1, &status, 0);
     if (done < 0) {
@@ -278,15 +297,17 @@ int runCoordinator() {
       continue;
     }
 
-    const int rank = (done == child0) ? 0 : 1;
+    int rank = 0;
+    while (rank < numRanks && children[rank] != done) { ++rank; }
+    if (rank == numRanks) { continue; }
     reaped[rank] = true;
     if (!WIFEXITED(status) || WEXITSTATUS(status) != kOk) {
       fprintf(stderr, "%s rank %d failed or terminated abnormally\n", tag,
               rank);
       if (!rc) { rc = kWorkerFailed; }
-      // Only signal the peer if it hasn't already been reaped
-      const int peer = 1 - rank;
-      if (!reaped[peer]) { kill(children[peer], SIGKILL); }
+      for (int peer = 0; peer < numRanks; ++peer) {
+        if (!reaped[peer]) { kill(children[peer], SIGKILL); }
+      }
     }
   }
   return rc;
@@ -299,13 +320,14 @@ int runCoordinator() {
 TEST(AsymmetricVisibilityWorker, Run) {
   const char* rankEnv = getenv(kEnvRank);
   if (rankEnv == nullptr) {
-    GTEST_SKIP() << "worker entry point, driven by "
-                    "AsymmetricVisibility.CommInitRankAllReduce";
+    GTEST_SKIP() << "worker entry point, driven by AsymmetricVisibility.*";
   }
 
   const int rank = atoi(rankEnv);
+  const char* numRanksEnv = getenv(kEnvNumRanks);
   const char* devEnv = getenv(kEnvLocalDev);
   const char* uidEnv = getenv(kEnvUid);
+  ASSERT_NE(numRanksEnv, nullptr);
   ASSERT_NE(devEnv, nullptr);
   ASSERT_NE(uidEnv, nullptr);
 
@@ -314,15 +336,14 @@ TEST(AsymmetricVisibilityWorker, Run) {
 
   // Exit with the WorkerStatus code so the orchestrator's WEXITSTATUS reflects
   // this worker's own result rather than gtest's generic exit code.
-  const int rc = runWorker(rank, atoi(devEnv), id);
+  const int rc = runWorker(rank, atoi(numRanksEnv), atoi(devEnv), id);
   if (rc != kOk) { _exit(rc); }
 }
 
-// Orchestrator: reproduces the ROCM-27034 asymmetric topology with two workers,
-// rank0 seeing devices {0,1} and binding ordinal 1, rank1 seeing device {2} and
-// binding ordinal 0. All HIP/NCCL work happens in a forked coordinator so the
-// main gtest process stays HIP-clean for TestBed's per-test forks.
-TEST(AsymmetricVisibility, CommInitRankAllReduce) {
+// Runs the workers of a layout. All HIP/NCCL work happens in a forked
+// coordinator so the main gtest process stays HIP-clean for TestBed's per-test
+// forks.
+static void runLayout(VisibilityLayout layout) {
   const char* cumemEnv = getenv("NCCL_CUMEM_ENABLE");
   const int cumemValue = cumemEnv != nullptr ? atoi(cumemEnv) : 1;
   if (cumemValue == 0) {
@@ -331,7 +352,7 @@ TEST(AsymmetricVisibility, CommInitRankAllReduce) {
   }
 
   const pid_t coord = fork();
-  if (coord == 0) { _exit(runCoordinator()); }
+  if (coord == 0) { _exit(runCoordinator(layout)); }
   ASSERT_GT(coord, 0) << "fork of coordinator process failed";
 
   int status = 0;
@@ -339,12 +360,15 @@ TEST(AsymmetricVisibility, CommInitRankAllReduce) {
   ASSERT_TRUE(WIFEXITED(status)) << "coordinator terminated abnormally";
   const int code = WEXITSTATUS(status);
   if (code == kCoordSkip) {
-    GTEST_SKIP() << "requires at least " << kMinDevices
-                 << " GPUs for an asymmetric topology";
+    GTEST_SKIP() << "not enough GPUs for this visibility layout";
   }
   EXPECT_EQ(code, kOk)
     << "asymmetric-visibility workers reported failure (exit code " << code
     << "); see stderr for the failing rank";
 }
 
+TEST(AsymmetricVisibility, CommInitRankAllReduce) { runLayout(asymmetricLayout); }
+
+TEST(AsymmetricVisibility, OneGpuPerRankAllReduce) { runLayout(oneGpuPerRankLayout); }
+
 } // namespace RcclUnitTesting
diff --git a/projects/rccl/test/graph/TopoGpuDevTests.cpp b/projects/rccl/test/graph/TopoGpuDevTests.cpp
index 43d81dba4b7c..5d49b7e293fc 100644
--- a/projects/rccl/test/graph/TopoGpuDevTests.cpp
+++ b/projects/rccl/test/graph/TopoGpuDevTests.cpp
@@ -5,7 +5,8 @@
  ************************************************************************/
 
 // Checks that the GPU nodes of a real topology carry the SMI index as `dev`, distinct per GPU.
-// Needs at least two GPUs; on a host whose HIP order matches the SMI order it cannot fail.
+// A check can only fail where HIP order differs from SMI order, so it runs once with the host's
+// order and once with the first two GPUs swapped; one of the two orders differs on any host.
 
 #include <gtest/gtest.h>
 #include <rccl/rccl.h>
@@ -19,32 +20,45 @@
 
 namespace RcclUnitTesting {
 
-TEST(TopoGpuDev, OwnGpuDevIsNvmlDevAndUnique)
-{
-  RUN_ISOLATED_TEST("TopoGpuDev_OwnGpuDevIsNvmlDevAndUnique", []() {
-    int numDevices = 0;
-    ASSERT_EQ(hipGetDeviceCount(&numDevices), hipSuccess);
-    if (numDevices < 2) GTEST_SKIP() << "Needs at least 2 GPUs";
-
-    std::vector<ncclComm_t> comms(numDevices, nullptr);
-    ASSERT_EQ(ncclCommInitAll(comms.data(), numDevices, nullptr), ncclSuccess);
-
-    for (ncclComm_t comm : comms) {
-      const struct ncclTopoNodeSet& gpus = comm->topo->nodes[GPU];
-      int ownDev = -1;
-      for (int gpuIndex = 0; gpuIndex < gpus.count; gpuIndex++) {
-        if (gpus.nodes[gpuIndex].gpu.rank == comm->rank) ownDev = gpus.nodes[gpuIndex].gpu.dev;
-      }
-      EXPECT_EQ(ownDev, comm->nvmlDev) << "rank " << comm->rank << " (HIP device " << comm->cudaDev << ")";
+namespace {
+
+void checkGpuDevsAreNvmlDevs() {
+  int numDevices = 0;
+  ASSERT_EQ(hipGetDeviceCount(&numDevices), hipSuccess);
+  if (numDevices < 2) GTEST_SKIP() << "Needs at least 2 GPUs";
+
+  std::vector<ncclComm_t> comms(numDevices, nullptr);
+  ASSERT_EQ(ncclCommInitAll(comms.data(), numDevices, nullptr), ncclSuccess);
+
+  for (ncclComm_t comm : comms) {
+    const struct ncclTopoNodeSet& gpus = comm->topo->nodes[GPU];
+    int ownDev = -1;
+    for (int gpuIndex = 0; gpuIndex < gpus.count; gpuIndex++) {
+      if (gpus.nodes[gpuIndex].gpu.rank == comm->rank) ownDev = gpus.nodes[gpuIndex].gpu.dev;
     }
+    EXPECT_EQ(ownDev, comm->nvmlDev) << "rank " << comm->rank << " (HIP device " << comm->cudaDev << ")";
+  }
+
+  const struct ncclTopoNodeSet& gpus = comms[0]->topo->nodes[GPU];
+  std::set<int> devs;
+  for (int gpuIndex = 0; gpuIndex < gpus.count; gpuIndex++) devs.insert(gpus.nodes[gpuIndex].gpu.dev);
+  EXPECT_EQ(static_cast<int>(devs.size()), gpus.count) << "GPU nodes share a dev index";
 
-    const struct ncclTopoNodeSet& gpus = comms[0]->topo->nodes[GPU];
-    std::set<int> devs;
-    for (int gpuIndex = 0; gpuIndex < gpus.count; gpuIndex++) devs.insert(gpus.nodes[gpuIndex].gpu.dev);
-    EXPECT_EQ(static_cast<int>(devs.size()), gpus.count) << "GPU nodes share a dev index";
+  for (ncclComm_t comm : comms) EXPECT_EQ(ncclCommDestroy(comm), ncclSuccess);
+}
 
-    for (ncclComm_t comm : comms) EXPECT_EQ(ncclCommDestroy(comm), ncclSuccess);
-  });
+} // namespace
+
+TEST(TopoGpuDev, OwnGpuDevIsNvmlDevAndUnique)
+{
+  RUN_ISOLATED_TEST("TopoGpuDev_OwnGpuDevIsNvmlDevAndUnique", []() { checkGpuDevsAreNvmlDevs(); });
+}
+
+TEST(TopoGpuDev, OwnGpuDevIsNvmlDevWithSwappedVisibility)
+{
+  RUN_ISOLATED_TEST_WITH_ENV(
+    "TopoGpuDev_OwnGpuDevIsNvmlDevWithSwappedVisibility", []() { checkGpuDevsAreNvmlDevs(); },
+    {{"HIP_VISIBLE_DEVICES", "1,0"}});
 }
 
 } // namespace RcclUnitTesting
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json
index 47fe2814f84c..fea0ec7e58e1 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json
@@ -1346,7 +1346,8 @@
         "NCCL_DEBUG": "INFO"
       },
       "tests": [
-        {"name": "AsymmetricVisibility.CommInitRankAllReduce", "description": "ROCM-27034: RCCL comm init + AllReduce under asymmetric HIP_VISIBLE_DEVICES", "test_filter": "AsymmetricVisibility.CommInitRankAllReduce"}
+        {"name": "AsymmetricVisibility.CommInitRankAllReduce", "description": "ROCM-27034: RCCL comm init + AllReduce under asymmetric HIP_VISIBLE_DEVICES", "test_filter": "AsymmetricVisibility.CommInitRankAllReduce"},
+        {"name": "AsymmetricVisibility.OneGpuPerRankAllReduce", "description": "Topology GPU dev index when each rank sees only its own GPU", "test_filter": "AsymmetricVisibility.OneGpuPerRankAllReduce"}
       ]
     },
     "plugin_compat_v10": {
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json
index 93b6458d08ed..d6c347dfd306 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json
@@ -1533,6 +1533,11 @@
           "name": "AsymmetricVisibility.CommInitRankAllReduce",
           "description": "ROCM-27034: RCCL comm init + AllReduce under asymmetric HIP_VISIBLE_DEVICES",
           "test_filter": "AsymmetricVisibility.CommInitRankAllReduce"
+        },
+        {
+          "name": "AsymmetricVisibility.OneGpuPerRankAllReduce",
+          "description": "Topology GPU dev index when each rank sees only its own GPU",
+          "test_filter": "AsymmetricVisibility.OneGpuPerRankAllReduce"
         }
       ]
     },
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json
index 83959e389a6f..f7a68fe5a0a0 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json
@@ -4040,6 +4040,11 @@
           "name": "AsymmetricVisibility.CommInitRankAllReduce",
           "description": "ROCM-27034: RCCL comm init + AllReduce under asymmetric HIP_VISIBLE_DEVICES",
           "test_filter": "AsymmetricVisibility.CommInitRankAllReduce"
+        },
+        {
+          "name": "AsymmetricVisibility.OneGpuPerRankAllReduce",
+          "description": "Topology GPU dev index when each rank sees only its own GPU",
+          "test_filter": "AsymmetricVisibility.OneGpuPerRankAllReduce"
         }
       ]
     },
diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json
index e24d456638e8..ddd049733dad 100644
--- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json
+++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json
@@ -4333,6 +4333,11 @@
           "name": "AsymmetricVisibility.CommInitRankAllReduce",
           "description": "ROCM-27034: RCCL comm init + AllReduce under asymmetric HIP_VISIBLE_DEVICES",
           "test_filter": "AsymmetricVisibility.CommInitRankAllReduce"
+        },
+        {
+          "name": "AsymmetricVisibility.OneGpuPerRankAllReduce",
+          "description": "Topology GPU dev index when each rank sees only its own GPU",
+          "test_filter": "AsymmetricVisibility.OneGpuPerRankAllReduce"
         }
       ]
     },

From 83c921da6776aa707f7fb40bd7ebca6c84675607 Mon Sep 17 00:00:00 2001
From: Anatolii Rozanov <arozanov@amd.com>
Date: Thu, 15 Oct 2026 16:44:42 +0200
Subject: [PATCH 6/7] docs(rccl): drop the regression figure from the CHANGELOG
 entry

The measured slowdown depends on the system and the run, so the entry
names the gather performance regression without a number.

Co-authored-by: Cursor <cursoragent@cursor.com>
---
 projects/rccl/CHANGELOG.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md
index 365032cb638c..a5c0147e453d 100644
--- a/projects/rccl/CHANGELOG.md
+++ b/projects/rccl/CHANGELOG.md
@@ -48,7 +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, raising 16 KB to 1 MB `ncclGather` latency by up to 15% on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU.
+* 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.

From 3dac7a332ceb66cf3fef7051fb69326fc336e17e Mon Sep 17 00:00:00 2001
From: Anatolii Rozanov <arozanov@amd.com>
Date: Fri, 16 Oct 2026 11:09:14 +0200
Subject: [PATCH 7/7] test(rccl): exit asymmetric-visibility workers via _exit
 on success

On the coverage CI node the OneGpuPerRankAllReduce workers all passed
their gtest body but never exited, so the coordinator sat in waitpid
until the 120 s group timeout killed the run. A successful worker
returned through gtest and normal process teardown; only the failure
path called _exit.

Move the flush-stdio, write-coverage, _exit sequence of the re-exec
runner into ProcessIsolatedTestRunner::exitChild and have the workers
leave through it on every result, so their coverage is still written.

Co-authored-by: Cursor <cursoragent@cursor.com>
---
 .../rccl/test/AsymmetricVisibilityTests.cpp   |  4 +++-
 .../test/common/ProcessIsolatedTestRunner.cpp | 24 +++++++++++--------
 .../test/common/ProcessIsolatedTestRunner.hpp |  7 ++++++
 3 files changed, 24 insertions(+), 11 deletions(-)

diff --git a/projects/rccl/test/AsymmetricVisibilityTests.cpp b/projects/rccl/test/AsymmetricVisibilityTests.cpp
index b2891d6b917a..21d300f07621 100644
--- a/projects/rccl/test/AsymmetricVisibilityTests.cpp
+++ b/projects/rccl/test/AsymmetricVisibilityTests.cpp
@@ -40,6 +40,8 @@
 #include <string>
 #include <vector>
 
+#include "ProcessIsolatedTestRunner.hpp"
+
 namespace RcclUnitTesting {
 namespace {
 
@@ -337,7 +339,7 @@ TEST(AsymmetricVisibilityWorker, Run) {
   // Exit with the WorkerStatus code so the orchestrator's WEXITSTATUS reflects
   // this worker's own result rather than gtest's generic exit code.
   const int rc = runWorker(rank, atoi(numRanksEnv), atoi(devEnv), id);
-  if (rc != kOk) { _exit(rc); }
+  ProcessIsolatedTestRunner::exitChild(rc);
 }
 
 // Runs the workers of a layout. All HIP/NCCL work happens in a forked
diff --git a/projects/rccl/test/common/ProcessIsolatedTestRunner.cpp b/projects/rccl/test/common/ProcessIsolatedTestRunner.cpp
index d25659aa6b99..4d2fbd8db2a9 100644
--- a/projects/rccl/test/common/ProcessIsolatedTestRunner.cpp
+++ b/projects/rccl/test/common/ProcessIsolatedTestRunner.cpp
@@ -587,6 +587,19 @@ void ProcessIsolatedTestRunner::displayCapturedOutput(
     flush(std::cerr, output.stderrContent);
 }
 
+void ProcessIsolatedTestRunner::exitChild(int code)
+{
+    fflush(NULL);
+#if defined(RCCL_TEST_CODE_COVERAGE)
+    __llvm_profile_write_file();
+    using WriteFn = int (*)(void);
+    auto libWrite = reinterpret_cast<WriteFn>(
+        dlsym(RTLD_DEFAULT, "rcclCoverageWriteFile"));
+    if(libWrite) libWrite();
+#endif
+    _exit(code);
+}
+
 // Re-exec child entrypoint: if kReexecMarkerEnvVar is set, runs the named test
 // and _exit()s with the result. Returns normally only in the original parent.
 void ProcessIsolatedTestRunner::handleReexecEntrypoint(
@@ -603,16 +616,7 @@ void ProcessIsolatedTestRunner::handleReexecEntrypoint(
     {
         if(testConfig.name == target)
         {
-            int result = runTestInProcess(testConfig);
-            fflush(NULL);
-#if defined(RCCL_TEST_CODE_COVERAGE)
-            __llvm_profile_write_file();
-            using WriteFn = int (*)(void);
-            auto libWrite = reinterpret_cast<WriteFn>(
-                dlsym(RTLD_DEFAULT, "rcclCoverageWriteFile"));
-            if(libWrite) libWrite();
-#endif
-            _exit(result);
+            exitChild(runTestInProcess(testConfig));
         }
     }
 
diff --git a/projects/rccl/test/common/ProcessIsolatedTestRunner.hpp b/projects/rccl/test/common/ProcessIsolatedTestRunner.hpp
index 377d1338a94d..809b42911b87 100644
--- a/projects/rccl/test/common/ProcessIsolatedTestRunner.hpp
+++ b/projects/rccl/test/common/ProcessIsolatedTestRunner.hpp
@@ -325,6 +325,13 @@ class ProcessIsolatedTestRunner
      * @return Number of registered tests
      */
     static size_t getTestCount();
+
+    /**
+     * @brief Flush stdio and coverage data, then _exit() with the given code
+     * @note Use in forked or re-exec'd children: _exit() skips the atexit
+     *       handlers that would otherwise write coverage data.
+     */
+    [[noreturn]] static void exitChild(int code);
 };
 
 // Macros for Simplified Usage