diff --git a/projects/rccl/src/graph/topo.cc b/projects/rccl/src/graph/topo.cc
index aa7f08ad4fd9..8061f02dae9f 100644
--- a/projects/rccl/src/graph/topo.cc
+++ b/projects/rccl/src/graph/topo.cc
@@ -831,24 +831,13 @@ ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* s
   return ncclSuccess;
 }
 
-static bool ncclTopoXmlIsPrimaryGpuForDev(const struct ncclXmlNode* xmlGpu) {
-  // Only the first GPU entry (rank) for a given device (physical GPU) should contribute links.
-  for (int s = 0; s < xmlGpu->parent->nSubs; s++) {
-    const struct ncclXmlNode* sib = xmlGpu->parent->subs[s];
-    if (sib == xmlGpu) break;
-    if (strcmp(sib->name, "gpu") == 0) return false;
-  }
-  return true;
-}
-
 static ncclResult_t ncclTopoGetGpuDevNode(struct ncclXmlNode* xmlGpu, const char* busId, int systemId,
                                           struct ncclTopoSystem* system, struct ncclTopoNode** devNode) {
-  int mloPart, sm;
+  int mloPart;
   int64_t rawBusId;
   NCCLCHECK(busIdToInt64(busId, &rawBusId));
   NCCLCHECK(ncclTopoCheckMloPartBusId(rawBusId));
   NCCLCHECK(xmlGetAttrIntDefault(xmlGpu, "mlopart", &mloPart, NCCL_TOPO_UNDEF));
-  NCCLCHECK(xmlGetAttrInt(xmlGpu, "sm", &sm));
   int64_t devBusId =
     (mloPart != NCCL_TOPO_UNDEF && ncclParamTopoSplitMlopart()) ? NCCL_TOPO_MLOPART_BUSID(rawBusId, mloPart) : rawBusId;
   NCCLCHECK(ncclTopoGetNode(system, devNode, DEV, NCCL_TOPO_ID(systemId, devBusId)));
@@ -875,6 +864,29 @@ static ncclResult_t ncclTopoXmlIsPrimaryGpuForDev(struct ncclXmlNode* xmlGpu, bo
 }
 
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+static ncclResult_t ncclTopoGetDevForPciBusId(struct ncclTopoSystem* system, int systemId, int64_t busId,
+                                              struct ncclTopoNode** devNode) {
+  *devNode = NULL;
+  NCCLCHECK(ncclTopoGetNode(system, devNode, DEV, NCCL_TOPO_ID(systemId, busId)));
+  if (*devNode) return ncclSuccess;
+  // HIP uses the function nibble as the partition index. Match an overlaid DEV
+  // on the same PCI domain:bus:device, regardless of its physical function.
+  int fn = (int)(busId & 0xf);
+  if (fn >= NCCL_TOPO_MLOPART_DEV_MAX) return ncclSuccess;
+  int64_t phys = busId & ~0xfLL;
+  for (int d = 0; d < system->nodes[DEV].count; d++) {
+    struct ncclTopoNode* candidate = system->nodes[DEV].nodes + d;
+    int64_t local = NCCL_TOPO_ID_LOCAL_ID(candidate->id);
+    if (NCCL_TOPO_ID_SYSTEM_ID(candidate->id) == systemId &&
+        (local & ~(NCCL_TOPO_MLOPART_MASK | 0xfLL)) == phys &&
+        (local & NCCL_TOPO_MLOPART_MASK) == NCCL_TOPO_MLOPART(fn)) {
+      *devNode = candidate;
+      break;
+    }
+  }
+  return ncclSuccess;
+}
+
 // RCCL: XGMI is the AMD analogue of NVLink. Upstream v2.30 introduced the generic DEV
 // (physical device) topology node, with GPU (rank) nodes hanging off their parent DEV node
 // via LINK_LOC. ncclTopoComputePaths()/search now assemble inter-GPU links as DEV->DEV hops
@@ -886,14 +898,19 @@ static ncclResult_t ncclTopoXmlIsPrimaryGpuForDev(struct ncclXmlNode* xmlGpu, bo
 ncclResult_t ncclTopoAddXGMI(struct ncclXmlNode* node, struct ncclTopoSystem* system, const char* parentBusId,
                              int systemId) {
   if (strcmp(node->name, "xgmi") == 0) {
-    if (!ncclTopoXmlIsPrimaryGpuForDev(node->parent)) return ncclSuccess;
     struct ncclTopoNode* devNode = NULL;
-    int64_t pBusId;
-    NCCLCHECK(busIdToInt64(parentBusId, &pBusId));
-    pBusId = NCCL_TOPO_ID(systemId, pBusId);
-    NCCLCHECK(ncclTopoGetNode(system, &devNode, DEV, pBusId));
+    if (node->parent && strcmp(node->parent->name, "gpu") == 0) {
+      bool isPrimary = false;
+      NCCLCHECK(ncclTopoXmlIsPrimaryGpuForDev(node->parent, &isPrimary));
+      if (!isPrimary) return ncclSuccess;
+      NCCLCHECK(ncclTopoGetGpuDevNode(node->parent, parentBusId, systemId, system, &devNode));
+    } else {
+      int64_t pBusId;
+      NCCLCHECK(busIdToInt64(parentBusId, &pBusId));
+      NCCLCHECK(ncclTopoGetNode(system, &devNode, DEV, NCCL_TOPO_ID(systemId, pBusId)));
+    }
     if (devNode == NULL) {
-      WARN("Add XGMI error : could not find DEV for GPU %lx\n", pBusId);
+      WARN("Add XGMI error : could not find DEV for GPU %s\n", parentBusId);
       return ncclInternalError;
     }
     int count;
@@ -909,7 +926,7 @@ ncclResult_t ncclTopoAddXGMI(struct ncclXmlNode* node, struct ncclTopoSystem* sy
       NCCLCHECK(xmlGetAttrStr(node, "target", &target));
       int64_t busId;
       NCCLCHECK(busIdToInt64(target, &busId));
-      NCCLCHECK(ncclTopoGetNode(system, &remote, DEV, NCCL_TOPO_ID(systemId, busId)));
+      NCCLCHECK(ncclTopoGetDevForPciBusId(system, systemId, busId, &remote));
     } else if (targetType == CPU) {
       // XGMI connection to the local CPU
       NCCLCHECK(findLocalCpu(devNode, &remote, NULL));
@@ -1060,7 +1077,7 @@ ncclResult_t ncclTopoAddC2c(struct ncclXmlNode* node, struct ncclTopoSystem* sys
 
     if (nSibDevs > 1) {
       // Use a C2C bridge node to guarantee the total bw of the C2C link is shared between the devices.
-      // Note: pBusId is the dev busId; we have checked that the 2 last bits are 0 in ncclTopoAddGpuSub
+      // Note: pBusId is the dev busId; we have checked that the mlopart bits are 0 in ncclTopoAddGpuSub
       int64_t xc2cId = devNode->id & ~(int64_t)NCCL_TOPO_MLOPART_MASK;
       struct ncclTopoNode* xc2cNode = NULL;
       NCCLCHECK(ncclTopoGetNode(system, &xc2cNode, CXB, xc2cId));
diff --git a/projects/rccl/src/graph/topo.h b/projects/rccl/src/graph/topo.h
index 997cd130f02e..7a750bde8bb5 100644
--- a/projects/rccl/src/graph/topo.h
+++ b/projects/rccl/src/graph/topo.h
@@ -116,11 +116,21 @@ struct ncclTopoLinkList {
 #define RCCL_ROME_TOPO_PRESET_MODEL_IDX_4H4P (1000001)
 
 #define GCN_ARCH_NAME_LEN 16
-#define NCCL_TOPO_MLOPART_MASK (0x3) // lower 2 bits: bit[0]=enabled, bit[1]=partition index
-#define NCCL_TOPO_MLOPART_DEV_MAX (2) // max DEV nodes per physical GPU (one per uGPU partition)
-#define NCCL_TOPO_MLOPART(mloPart) ((((int64_t)(mloPart) << 1) | 0x1) & NCCL_TOPO_MLOPART_MASK)
+// The MLOPart partition index is overlaid on the DEV node's busId. These bits must sit above
+// the 36-bit PCI busId (domain[35:20] bus[19:12] device[11:4] function[3:0], see int64ToBusId)
+// and below NCCL_TOPO_GPU_LOCAL_RANK_SHIFT. AMD compute partitions expose HIP logical GPUs as
+// functions of the physical device: DPX/XCP use .0/.1, CPX uses .0 through .7. Only function
+// .0 exists in sysfs as a GPU; the rest are HIP aliases and share that PCI node via mlopart.
+#define NCCL_TOPO_MLOPART_SHIFT (36)
+#define NCCL_TOPO_MLOPART_DEV_MAX (8)
+#define NCCL_TOPO_MLOPART_MASK (((int64_t)0xf) << NCCL_TOPO_MLOPART_SHIFT) // enable bit + 3-bit index
+#define NCCL_TOPO_MLOPART(mloPart) \
+  (((((int64_t)(mloPart) << 1) | 0x1) << NCCL_TOPO_MLOPART_SHIFT) & NCCL_TOPO_MLOPART_MASK)
 #define NCCL_TOPO_MLOPART_BUSID(busId, mloPart) \
   ((mloPart) != NCCL_TOPO_UNDEF ? ((busId) | NCCL_TOPO_MLOPART(mloPart)) : (busId))
+static_assert(NCCL_TOPO_MLOPART_SHIFT >= 36, "MLOPart bits must sit above the 36-bit PCI busId");
+static_assert(NCCL_TOPO_MLOPART_SHIFT + 4 <= NCCL_TOPO_GPU_LOCAL_RANK_SHIFT,
+              "MLOPart bits overlap GPU local rank");
 
 struct ncclTopoNode {
   int type;
diff --git a/projects/rccl/src/graph/xml.cc b/projects/rccl/src/graph/xml.cc
index 12b29263755a..6084680a8b81 100644
--- a/projects/rccl/src/graph/xml.cc
+++ b/projects/rccl/src/graph/xml.cc
@@ -17,6 +17,7 @@
 #endif
 #include "archinfo.h"
 #include "os.h"
+#include "topo.h"
 #if defined(__x86_64__) || defined(_M_X64) || defined(_M_AMD64)
 #if defined(_MSC_VER)
 #include <intrin.h>
@@ -28,6 +29,33 @@
 // Arbitrarily large number for constructing virtual topology string
 #define NCCL_MAX_XML_DEPTH 1024
 
+static int ncclTopoPciClassIsGpu(const char* deviceClass) {
+  if (deviceClass == NULL || deviceClass[0] == '\0') return 0;
+  if (strncmp(deviceClass, PCI_ACCELERATOR_CLASS, strlen(PCI_ACCELERATOR_CLASS)) == 0) return 1;
+  if (strncmp(deviceClass, "0x03", 4) == 0) return 1;
+  return 0;
+}
+
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+// XGMI peers come from accelerator enumeration, so do not infer their class
+// from a HIP alias BDF that may be missing or identify another PCI function.
+static ncclResult_t ncclTopoXmlAddXgmiPeerGpu(struct ncclXml* xml, struct ncclXmlNode* gpuNode, const char* target,
+                                              int count, struct ncclXmlNode** nodeOut, int* createdOut = NULL) {
+  struct ncclXmlNode* node = NULL;
+  NCCLCHECK(xmlGetSubKv(gpuNode, "xgmi", &node, "target", target));
+  int created = node == NULL;
+  if (created) {
+    NCCLCHECK(xmlAddNode(xml, gpuNode, "xgmi", &node));
+    NCCLCHECK(xmlSetAttr(node, "target", target));
+    NCCLCHECK(xmlSetAttrInt(node, "count", count));
+    NCCLCHECK(xmlSetAttr(node, "tclass", PCI_ACCELERATOR_CLASS));
+  }
+  if (nodeOut) *nodeOut = node;
+  if (createdOut) *createdOut = created;
+  return ncclSuccess;
+}
+#endif
+
 /*******************/
 /* XML File Parser */
 /*******************/
@@ -366,7 +394,7 @@ ncclResult_t ncclTopoXmlLoadGpu(FILE* file, struct ncclXml* xml, struct ncclXmlN
 #else
   struct xmlHandler handlers[] = {{"nvlink", ncclTopoXmlLoadNvlink}, {"c2c", ncclTopoXmlLoadC2c}};
 #endif
-  NCCLCHECK(xmlLoadSub(file, xml, head, handlers, 2));
+  NCCLCHECK(xmlLoadSub(file, xml, head, handlers, sizeof(handlers) / sizeof(handlers[0])));
   return ncclSuccess;
 }
 
@@ -905,10 +933,48 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml*
   return ret;
 }
 
-ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev, struct ncclXml* xml,
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+// HIP may expose logical GPUs at function .1-.7 of the same device while sysfs
+// 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) {
+  const char* pciBusId = NULL;
+  NCCLCHECK(xmlGetAttr(pciNode, "busid", &pciBusId));
+  if (pciBusId == NULL) return ncclSuccess;
+  int64_t physId = 0;
+  NCCLCHECK(busIdToInt64(pciBusId, &physId));
+  physId &= ~0xfLL;
+  int nHip = 0;
+  if (hipGetDeviceCount(&nHip) != hipSuccess) return ncclSuccess;
+  for (int h = 0; h < nHip; h++) {
+    if (h == (int)xmlDev) continue;
+    char peerBus[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];
+    if (hipDeviceGetPCIBusId(peerBus, sizeof(peerBus), h) != hipSuccess) continue;
+    int64_t peerId = 0;
+    if (busIdToInt64(peerBus, &peerId) != ncclSuccess) continue;
+    if ((peerId & ~0xfLL) != physId) continue;
+    char lowerId[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];
+    for (int c = 0; c < NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE; c++) {
+      lowerId[c] = tolower(peerBus[c]);
+      if (peerBus[c] == 0) break;
+    }
+    struct ncclXmlNode* nvl = NULL;
+    NCCLCHECK(ncclTopoXmlAddXgmiPeerGpu(xml, gpuNode, lowerId, 1, &nvl));
+    if (nvl) {
+      int idx;
+      NCCLCHECK(xmlGetAttrIndex(nvl, "fabric_supported", &idx));
+      if (idx == -1) NCCLCHECK(xmlSetAttrInt(nvl, "fabric_supported", 0));
+    }
+  }
+  return ncclSuccess;
+}
+#endif
+
+ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t xmlDev, uint32_t smiDev, struct ncclXml* xml,
                                    struct ncclXmlNode** gpuNodeRet) {
   struct ncclXmlNode* gpuNode = NULL;
-  NCCLCHECK(xmlGetSub(pciNode, "gpu", &gpuNode));
+  NCCLCHECK(xmlGetSubKvInt(pciNode, "gpu", &gpuNode, "dev", (int)xmlDev));
   if (gpuNode == NULL) NCCLCHECK(xmlAddNode(xml, pciNode, "gpu", &gpuNode));
 
   int index = -1;
@@ -916,7 +982,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
   int dev = -1;
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "dev", &index));
   if (index == -1) {
-    NCCLCHECK(xmlSetAttrInt(gpuNode, "dev", rocmDev));
+    NCCLCHECK(xmlSetAttrInt(gpuNode, "dev", xmlDev));
   }
   NCCLCHECK(xmlGetAttrInt(gpuNode, "dev", &dev));
   if (dev == -1) {
@@ -927,7 +993,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "sm", &index));
   if (index == -1) {
     cudaDeviceProp devProp;
-    CUDACHECK(cudaGetDeviceProperties(&devProp, 0));
+    CUDACHECK(cudaGetDeviceProperties(&devProp, (int)xmlDev));
     NCCLCHECK(xmlSetAttrInt(gpuNode, "sm", devProp.multiProcessorCount));
   }
   int sm;
@@ -938,7 +1004,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "gcn", &index));
   if (index == -1) {
     hipDeviceProp_t devProp;
-    CUDACHECK(hipGetDeviceProperties(&devProp, 0));
+    CUDACHECK(hipGetDeviceProperties(&devProp, (int)xmlDev));
     // extract only the releveant info from the gcnArchName attribute
     // e.g.: convert "gfx908:sramecc+:xnack-" to "gfx908"
     char gcnArchNameSubstr[128];
@@ -956,7 +1022,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
   NCCLCHECK(xmlGetAttrIndex(gpuNode, "arch", &index));
   if (index == -1) {
     hipDeviceProp_t devProp;
-    CUDACHECK(hipGetDeviceProperties(&devProp, 0));
+    CUDACHECK(hipGetDeviceProperties(&devProp, (int)xmlDev));
     memcpy(&arch.arch, &devProp.arch, sizeof(hipDeviceArch_t));
     NCCLCHECK(xmlSetAttrInt(gpuNode, "arch", arch.value));
   }
@@ -976,14 +1042,14 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
 #ifdef USE_AMDSMI
     NCCLCHECK(amd_smi_getNumDevice(&deviceCnt));
     for (int i = 0; i < deviceCnt; i++) {
-      if (i != dev) {
+      if (i != smiDev) {
         amdsmi_link_type_t amdsmi_type;
         // Based on amd-smi team, there will be a new UALoE link type
         // First, we will check if UALoE is supported, if not we will check for XGMI links
         bool canUseUALoE = false;
         amdsmiFabricDeviceInfo fabInfo_dev;
         amdsmiFabricDeviceInfo fabInfo;
-        if (amd_smi_getFabricDeviceInfo(dev, &fabInfo_dev) == ncclSuccess &&
+        if (amd_smi_getFabricDeviceInfo(smiDev, &fabInfo_dev) == ncclSuccess &&
             amd_smi_getFabricDeviceInfo(i, &fabInfo) == ncclSuccess) {
           if (fabInfo_dev.fabricSupported && fabInfo.fabricSupported && fabInfo_dev.cliqueId == fabInfo.cliqueId &&
               (memcmp(fabInfo_dev.clusterUuid, fabInfo.clusterUuid, UALOE_GPU_FABRIC_UUID_LEN) == 0)) {
@@ -998,10 +1064,16 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
             lowerId[c] = tolower(busIdStr[c]);
             if (busIdStr[c] == 0) break;
           }
-          NCCLCHECK(xmlGetSubKv(gpuNode, "xgmi", &nvlNode, "target", lowerId));
-          if (nvlNode == NULL) {
-            NCCLCHECK(xmlAddNode(xml, gpuNode, "xgmi", &nvlNode));
-            NCCLCHECK(xmlSetAttr(nvlNode, "target", lowerId));
+          // UALoE when both peers report a usable fabric (ready/active, same clique).
+          // Unconfigured ualink is a node setup issue, not a missing interconnect:
+          // fabricSupported is false and discovery falls through to 1-hop XGMI.
+          // Either path stamps accelerator tclass: HIP bus IDs need not exist on
+          // the physical PCI bus, so sysfs class at that BDF is not the peer GPU.
+          int created;
+          // count=1: amd-smi does not yet report UALoE link count.
+          // TODO: Update amd_smi_getLinkInfo to return the number of links for UALoE.
+          NCCLCHECK(ncclTopoXmlAddXgmiPeerGpu(xml, gpuNode, lowerId, 1, &nvlNode, &created));
+          if (created) {
             uint64_t uuidHigh = 0;
             uint64_t uuidLow = 0;
             memcpy(&uuidHigh, fabInfo.clusterUuid, sizeof(uuidHigh));
@@ -1015,15 +1087,10 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
             NCCLCHECK(xmlSetAttrInt(nvlNode, "ppod_size", fabInfo.ppodSize));
             NCCLCHECK(xmlSetAttrInt(nvlNode, "vpod_id", fabInfo.cliqueId));
             NCCLCHECK(xmlSetAttrInt(nvlNode, "vpod_size", fabInfo.vpodSize));
-            // We don't have a way to get the number of links between two devices when using UALoE, so we will set it to 1 to avoid breaking topology construction.
-            // This is not ideal, but it will allow us to use UALoE without having to wait for a new version of amd-smi that returns the number of links.
-            // We can update this later when we have that information available.
-            // TODO: Update amd_smi_getLinkInfo to return the number of links for UALoE and update this code accordingly.
-            NCCLCHECK(xmlSetAttrInt(nvlNode, "count", 1));
           }
         } else {
           int hops, count;
-          if (amd_smi_getLinkInfo(dev, i, &amdsmi_type, &hops, &count) == ncclSuccess) {
+          if (amd_smi_getLinkInfo(smiDev, i, &amdsmi_type, &hops, &count) == ncclSuccess) {
             if (amdsmi_type == AMDSMI_LINK_TYPE_XGMI && hops == 1) {
               char busIdStr[] = "00000000:00:00.0";
               NCCLCHECK(amd_smi_getDevicePciBusIdString(i, busIdStr, sizeof(busIdStr)));
@@ -1032,11 +1099,9 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
                 lowerId[c] = tolower(busIdStr[c]);
                 if (busIdStr[c] == 0) break;
               }
-              NCCLCHECK(xmlGetSubKv(gpuNode, "xgmi", &nvlNode, "target", lowerId));
-              if (nvlNode == NULL) {
-                NCCLCHECK(xmlAddNode(xml, gpuNode, "xgmi", &nvlNode));
-                NCCLCHECK(xmlSetAttr(nvlNode, "target", lowerId));
-                NCCLCHECK(xmlSetAttrInt(nvlNode, "count", count));
+              int created;
+              NCCLCHECK(ncclTopoXmlAddXgmiPeerGpu(xml, gpuNode, lowerId, count, &nvlNode, &created));
+              if (created) {
                 NCCLCHECK(xmlSetAttrInt(nvlNode, "fabric_supported", 0));
               }
             }
@@ -1047,10 +1112,10 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
 #else
     NCCLCHECK(rocm_smi_getNumDevice(&deviceCnt));
     for (int i = 0; i < deviceCnt; i++) {
-      if (i != dev) {
+      if (i != smiDev) {
         RSMI_IO_LINK_TYPE rsmi_type;
         int hops, count;
-        if (rocm_smi_getLinkInfo(dev, i, &rsmi_type, &hops, &count) == ncclSuccess) {
+        if (rocm_smi_getLinkInfo(smiDev, i, &rsmi_type, &hops, &count) == ncclSuccess) {
           if (rsmi_type == RSMI_IOLINK_TYPE_XGMI && hops == 1) {
             char busIdStr[] = "00000000:00:00.0";
             NCCLCHECK(rocm_smi_getDevicePciBusIdString(i, busIdStr, sizeof(busIdStr)));
@@ -1059,12 +1124,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
               lowerId[c] = tolower(busIdStr[c]);
               if (busIdStr[c] == 0) break;
             }
-            NCCLCHECK(xmlGetSubKv(gpuNode, "xgmi", &nvlNode, "target", lowerId));
-            if (nvlNode == NULL) {
-              NCCLCHECK(xmlAddNode(xml, gpuNode, "xgmi", &nvlNode));
-              NCCLCHECK(xmlSetAttr(nvlNode, "target", lowerId));
-              NCCLCHECK(xmlSetAttrInt(nvlNode, "count", count));
-            }
+            NCCLCHECK(ncclTopoXmlAddXgmiPeerGpu(xml, gpuNode, lowerId, count, &nvlNode));
           }
         }
       }
@@ -1140,6 +1200,9 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
     }
 #endif
   }
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+  NCCLCHECK(ncclTopoXmlAddHipPartitionXgmi(xml, gpuNode, pciNode, xmlDev));
+#endif
 #if CUDART_VERSION >= 11080
   struct ncclXmlNode* c2cNode = NULL;
   NCCLCHECK(xmlGetSub(gpuNode, "c2c", &c2cNode));
@@ -1173,7 +1236,11 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
     }
   }
 #endif
-  // Fill target classes
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+  // Sysfs class is only meaningful for a physical PCI function. HIP can name a
+  // logical GPU at a BDF that is not on the bus (e.g. function .1); ualink
+  // configured or not does not change that. Trust only GPU/accelerator classes.
+#endif
   for (int s = 0; s < gpuNode->nSubs; s++) {
     struct ncclXmlNode* sub = gpuNode->subs[s];
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
@@ -1195,11 +1262,29 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
         deviceClass[0] = '\0';
         if (ncclOsGetPciDeviceClassByBusId(busId, deviceClass, sizeof(deviceClass)) == ncclSuccess &&
             deviceClass[0] != '\0') {
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+          int isGpu = ncclTopoPciClassIsGpu(deviceClass);
+          if (!isGpu) {
+            INFO(NCCL_GRAPH,
+                 "XGMI target %s sysfs class %s is not a GPU (HIP logical BDF?); using accelerator tclass",
+                 busId, deviceClass);
+            NCCLCHECK(xmlSetAttr(sub, "tclass", PCI_ACCELERATOR_CLASS));
+          } else {
+            NCCLCHECK(xmlSetAttr(sub, "tclass", deviceClass));
+            TRACE(NCCL_GRAPH, "Read NVLink target class: tclass=%s for busId=%s", deviceClass, busId);
+          }
+#else
           NCCLCHECK(xmlSetAttr(sub, "tclass", deviceClass));
           TRACE(NCCL_GRAPH, "Read NVLink target class: tclass=%s for busId=%s", deviceClass, busId);
+#endif
         } else {
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+          INFO(NCCL_GRAPH, "No sysfs PCI class for XGMI target %s; using accelerator tclass", busId);
+          NCCLCHECK(xmlSetAttr(sub, "tclass", PCI_ACCELERATOR_CLASS));
+#else
           INFO(NCCL_GRAPH, "Could not get device class for NVLink target %s, assuming NVSwitch", busId);
           NCCLCHECK(xmlSetAttr(sub, "tclass", PCI_NVSWITCH_CLASS));
+#endif
         }
       }
     }
@@ -1210,25 +1295,50 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
 
 ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct ncclXmlNode** gpuNode) {
   struct ncclXmlNode* node;
-  NCCLCHECK(ncclTopoGetPciNode(xml, busId, &node));
+  const char* pciBusId = busId;
+  char physBusId[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+  // HIP logical GPUs can sit at a BDF that is not the GPU's PCI function (DPX
+  // function .1 is often a NIC in sysfs; CPX .1-.7 are missing from sysfs).
+  // Attach the <gpu> under the physical function-0 PCI node; mlopart is filled
+  // from the fake function in fillInfo.
+  int64_t hipId = 0;
+  NCCLCHECK(busIdToInt64(busId, &hipId));
+  int fn = (int)(hipId & 0xf);
+  if (fn > 0 && fn < NCCL_TOPO_MLOPART_DEV_MAX) {
+    char deviceClass[MAX_STR_LEN];
+    deviceClass[0] = '\0';
+    (void)ncclOsGetPciDeviceClassByBusId(busId, deviceClass, sizeof(deviceClass));
+    if (!ncclTopoPciClassIsGpu(deviceClass)) {
+      NCCLCHECK(int64ToBusId(hipId & ~0xfLL, physBusId));
+      pciBusId = physBusId;
+      INFO(NCCL_GRAPH,
+           "HIP logical GPU at %s is not a physical GPU PCI function (class %s); attaching under %s", busId,
+           deviceClass[0] ? deviceClass : "missing", pciBusId);
+    }
+  }
+#endif
+  NCCLCHECK(ncclTopoGetPciNode(xml, pciBusId, &node));
   NCCLCHECK(xmlSetAttrIfUnset(node, "class", "0x03"));
   NCCLCHECK(ncclTopoGetXmlFromSys(node, xml));
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
-  uint32_t devIndex = 0;
+  uint32_t smiIndex = 0;
 #ifdef USE_AMDSMI
   static int amdsmiInit = 0;
   if (amdsmiInit == 0) {
     NCCLCHECK(amd_smi_init());
   }
-  NCCLCHECK(amd_smi_getDeviceIndexByPciBusId(busId, &devIndex));
+  NCCLCHECK(amd_smi_getDeviceIndexByPciBusId(busId, &smiIndex));
 #else
   static int rocmsmiInit = 0;
   if (rocmsmiInit == 0) {
     NCCLCHECK(rocm_smi_init());
   }
-  NCCLCHECK(rocm_smi_getDeviceIndexByPciBusId(busId, &devIndex));
+  NCCLCHECK(rocm_smi_getDeviceIndexByPciBusId(busId, &smiIndex));
 #endif
-  NCCLCHECK(ncclTopoGetXmlFromGpu(node, devIndex, xml, gpuNode));
+  int hipDev = -1;
+  if (hipDeviceGetByPCIBusId(&hipDev, busId) != hipSuccess || hipDev < 0) hipDev = (int)smiIndex;
+  NCCLCHECK(ncclTopoGetXmlFromGpu(node, (uint32_t)hipDev, smiIndex, xml, gpuNode));
 #else
   nvmlDevice_t nvmlDev;
   NCCLCHECK(ncclNvmlDeviceGetHandleByPciBusId(busId, &nvmlDev));
diff --git a/projects/rccl/src/graph/xml.h b/projects/rccl/src/graph/xml.h
index ff92e51008ae..b8d1b3358948 100644
--- a/projects/rccl/src/graph/xml.h
+++ b/projects/rccl/src/graph/xml.h
@@ -19,6 +19,8 @@
 
 // PCI device class for NVSwitch (used to identify remote NVLink targets)
 #define PCI_NVSWITCH_CLASS "0x068000"
+// PCI device class for AMD accelerators ("Processing accelerators"), mapped to GPU in kvDictPciClass
+#define PCI_ACCELERATOR_CLASS "0x120000"
 
 // A few constraints to make the implementation easy
 #define MAX_STR_LEN 255
diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc
index 3fab7902b794..566ea6b2cdfc 100644
--- a/projects/rccl/src/init.cc
+++ b/projects/rccl/src/init.cc
@@ -1083,6 +1083,32 @@ static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, u
   info->shmDev = statbuf.st_dev;
 #endif
   info->busId = comm->busId;
+#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+  // HIP names do not contain "MLOPart". DPX/XCP/CPX logical GPUs are still
+  // exposed as PCI function .N while sysfs at that BDF is not a GPU (often the
+  // BDF is missing entirely). Use the PCI function as the partition index so
+  // ranks share the physical function-0 PCI node (see ncclTopoFillGpu) with
+  // distinct overlay DEV ids: function 0 is mlopart 0, a fake non-GPU function
+  // is mlopart N (CPX uses N=1..7).
+  if (info->mloPart == NCCL_TOPO_UNDEF) {
+    int fn = (int)(info->busId & 0xf);
+    if (fn < NCCL_TOPO_MLOPART_DEV_MAX) {
+      char busIdStr[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];
+      char deviceClass[MAX_STR_LEN];
+      deviceClass[0] = '\0';
+      if (int64ToBusId(info->busId, busIdStr) == ncclSuccess) {
+        (void)ncclOsGetPciDeviceClassByBusId(busIdStr, deviceClass, sizeof(deviceClass));
+        int isGpu = strncmp(deviceClass, PCI_ACCELERATOR_CLASS, strlen(PCI_ACCELERATOR_CLASS)) == 0 ||
+                    strncmp(deviceClass, "0x03", 4) == 0;
+        if (fn == 0) {
+          if (isGpu) info->mloPart = 0;
+        } else if (!isGpu) {
+          info->mloPart = fn;
+        }
+      }
+    }
+  }
+#endif
   CUCHECK(cuDeviceGetUuid((CUuuid*)&info->gpuUuid, (CUdevice)comm->cudaDev));
 
   // detect if fine grained memory is available on this GPU
@@ -1797,7 +1823,9 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
   // AllGather3 - begin
   NCCLCHECKGOTO(ncclCalloc(&allGather3Data, nranks), ret, fail);
   int idx;
-  NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, NCCL_TOPO_ID(comm->topo->systemId, comm->busId), &idx));
+  // GPU node ids include the MLOPart overlay (and a local-rank-on-DEV field), so they
+  // no longer match the raw PCI busId. Look up this rank's GPU node instead.
+  NCCLCHECK(ncclTopoRankToIndex(comm->topo, rank, &idx, /*showWarn=*/true));
   allGather3Data[rank].nc = 2;
   if (comm->topo->nodes[GPU].count == comm->topo->nRanks &&
       IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx906") && allXgmi)
diff --git a/projects/rccl/src/misc/alt_rsmi.cc b/projects/rccl/src/misc/alt_rsmi.cc
index 002d02ba2d87..bcf3081b2ca2 100644
--- a/projects/rccl/src/misc/alt_rsmi.cc
+++ b/projects/rccl/src/misc/alt_rsmi.cc
@@ -382,7 +382,11 @@ int ARSMI_get_fw_version(uint32_t /*dv_ind*/, uint64_t* fw_version) {
 
 // Find the DRM card number whose PCI slot name matches the given BDF.
 // Returns the card number on success, -1 if not found.
-static int findCardForBdf(uint32_t domain, uint8_t bus, uint8_t device, uint8_t function) {
+//
+// HIP/KFD may encode a DPX/XCP partition as PCI function N when the only
+// physical endpoint is function 0. If the exact BDF is missing, retry with
+// the function field stripped.
+static int findCardForBdfExact(uint32_t domain, uint8_t bus, uint8_t device, uint8_t function) {
   for (uint32_t card = 0; card < 128; card++) {
     char path[256];
     snprintf(path, sizeof(path), "%s/card%u/device/uevent", kDrmClassRoot, card);
@@ -404,6 +408,13 @@ static int findCardForBdf(uint32_t domain, uint8_t bus, uint8_t device, uint8_t
   return -1;
 }
 
+static int findCardForBdf(uint32_t domain, uint8_t bus, uint8_t device, uint8_t function) {
+  int card = findCardForBdfExact(domain, bus, device, function);
+  if (card >= 0) return card;
+  if (function != 0) return findCardForBdfExact(domain, bus, device, 0);
+  return -1;
+}
+
 // Read a sysfs attribute as a NUL-terminated, newline-stripped string.
 // Returns 0 on success, -1 if the file cannot be opened or is empty.
 static int readUalinkStr(const char* dir, const char* attr, char* buf, size_t len) {
diff --git a/projects/rccl/test/AltRsmiTests.cpp b/projects/rccl/test/AltRsmiTests.cpp
index 338ec8b4c66a..7c3cb143cbe7 100644
--- a/projects/rccl/test/AltRsmiTests.cpp
+++ b/projects/rccl/test/AltRsmiTests.cpp
@@ -1753,4 +1753,50 @@ TEST(AltRsmiTest, FabricInfo_AddrModeUnknown) {
   );
 }
 
+TEST(AltRsmiTest, FabricInfo_FakePciFunctionFallsBackToFn0) {
+  RUN_ISOLATED_TEST(
+    "FabricInfo_FakePciFunctionFallsBackToFn0",
+    []() {
+      setupTestEnvironment();
+      // KFD location_id 23553 = 0x5C01 → bus 0x5c, device 0, function 1.
+      // DPX/XCP HIP names this as a GPU BDF, but the only DRM PCI slot is .0.
+      createDirectory("2");
+      createFile("2/gpu_id", "8192\n");
+      createFile("2/properties",
+                 "unique_id 99\n"
+                 "location_id 23553\n"
+                 "domain 0\n"
+                 "vendor_id 4098\n");
+      createDirectory("2/io_links");
+      AltRsmiTestUtils::ResetState();
+      removeDrmSandbox();
+      AltRsmiTestUtils::SetDrmRoot(kTestDrmBasePath);
+      ASSERT_EQ(ARSMI_init(), 0);
+      setupDrmCard(0, "0000:5c:00.0");
+      setupUalink(0, "UALoE", "active", 3, 1000, 10, 1, 0, 2, "source-aliasing", "");
+
+      uint32_t n = 0;
+      ASSERT_EQ(ARSMI_get_num_devices(&n), 0);
+      uint32_t fakeFn = n;
+      for (uint32_t i = 0; i < n; i++) {
+        uint64_t bdf = 0;
+        ASSERT_EQ(ARSMI_dev_pci_id_get(i, &bdf), 0);
+        if ((bdf & 0x7) == 1) {
+          fakeFn = i;
+          break;
+        }
+      }
+      ASSERT_LT(fakeFn, n);
+
+      ARSMI_fabricInfo info;
+      int result = ARSMI_get_fabric_info(fakeFn, &info);
+      ASSERT_EQ(result, 0);
+      ASSERT_EQ(info.supported, 1);
+      ASSERT_EQ(info.fabric_type, ARSMI_FABRIC_TYPE_UALOE);
+      removeDrmSandbox();
+      cleanupTestEnvironment();
+    }
+  );
+}
+
 } // namespace RcclUnitTesting
diff --git a/projects/rccl/test/graph/TopoTests.cpp b/projects/rccl/test/graph/TopoTests.cpp
index 69632cd5412d..1bd238583efe 100644
--- a/projects/rccl/test/graph/TopoTests.cpp
+++ b/projects/rccl/test/graph/TopoTests.cpp
@@ -16,8 +16,7 @@
 #include <cstdlib>
 #include <cstring>
 
-// busIdToInt64 is an internal helper (declared in utils.h). Forward-declare it
-// here so the test can compute node ids exactly the way the builder does.
+// busIdToInt64 is an internal helper (declared in utils.h).
 ncclResult_t busIdToInt64(const char* busId, int64_t* id);
 
 // ncclTopoAddXGMI is a non-static internal symbol in graph/topo.cc that is not
@@ -96,13 +95,14 @@ class TopoTest : public ::testing::Test {
     return pci;
   }
 
-  // GPU-to-GPU link entry (tclass 0x03 -> GPU) connecting to a peer GPU.
+  // GPU-to-GPU link entry connecting to a peer GPU. Default tclass 0x03 (VGA);
+  // xml.cc now also writes PCI_ACCELERATOR_CLASS ("0x120000") for AMD GPUs.
   void addGpuLink(struct ncclXmlNode* holder, const char* targetBusId,
-                  int count) {
+                  int count, const char* tclass = "0x03") {
     struct ncclXmlNode* link = nullptr;
     EXPECT_EQ(xmlAddNode(xml, holder, "xgmi", &link), ncclSuccess);
     EXPECT_EQ(xmlSetAttr(link, "target", targetBusId), ncclSuccess);
-    EXPECT_EQ(xmlSetAttr(link, "tclass", "0x03"), ncclSuccess);
+    EXPECT_EQ(xmlSetAttr(link, "tclass", tclass), ncclSuccess);
     EXPECT_EQ(xmlSetAttrInt(link, "count", count), ncclSuccess);
   }
 
@@ -126,13 +126,21 @@ class TopoTest : public ::testing::Test {
   // GPU <pci> node (with nested <gpu>) under a CPU, returning the pci node so
   // callers can attach <xgmi> link entries below it.
   struct ncclXmlNode* addGpuPci(struct ncclXmlNode* parent, const char* busId,
-                                const char* gcn, int rank, int dev) {
+                                const char* gcn, int rank, int dev,
+                                int mloPart = NCCL_TOPO_UNDEF) {
     struct ncclXmlNode* pci = nullptr;
     EXPECT_EQ(xmlAddNode(xml, parent, "pci", &pci), ncclSuccess);
     EXPECT_EQ(xmlSetAttr(pci, "busid", busId), ncclSuccess);
     EXPECT_EQ(xmlSetAttr(pci, "class", "0x03"), ncclSuccess); // 0x03 -> GPU
     EXPECT_EQ(xmlSetAttrInt(pci, "link_width", 16), ncclSuccess);
     EXPECT_EQ(xmlSetAttr(pci, "link_speed", "16.0 GT/s PCIe"), ncclSuccess);
+    addGpuUnderPci(pci, gcn, rank, dev, mloPart);
+    return pci;
+  }
+
+  struct ncclXmlNode* addGpuUnderPci(struct ncclXmlNode* pci, const char* gcn,
+                                     int rank, int dev,
+                                     int mloPart = NCCL_TOPO_UNDEF) {
     struct ncclXmlNode* gpu = nullptr;
     EXPECT_EQ(xmlAddNode(xml, pci, "gpu", &gpu), ncclSuccess);
     EXPECT_EQ(xmlSetAttrInt(gpu, "sm", 304), ncclSuccess);
@@ -141,7 +149,10 @@ class TopoTest : public ::testing::Test {
     EXPECT_EQ(xmlSetAttrInt(gpu, "rank", rank), ncclSuccess);
     EXPECT_EQ(xmlSetAttrInt(gpu, "dev", dev), ncclSuccess);
     EXPECT_EQ(xmlSetAttrInt(gpu, "gdr", 1), ncclSuccess);
-    return pci;
+    if (mloPart != NCCL_TOPO_UNDEF) {
+      EXPECT_EQ(xmlSetAttrInt(gpu, "mlopart", mloPart), ncclSuccess);
+    }
+    return gpu;
   }
 
   static struct ncclTopoLink* findLink(struct ncclTopoNode* from,
@@ -159,6 +170,17 @@ class TopoTest : public ::testing::Test {
   struct ncclXmlNode* root = nullptr;
 };
 
+// Encoding lives in headers; these run on every platform.
+TEST_F(TopoTest, MloPartBusId_DoesNotClobberPciFunction) {
+  int64_t busFn1 = 0;
+  ASSERT_EQ(busIdToInt64("0000:01:00.1", &busFn1), ncclSuccess);
+  const int64_t enc0 = NCCL_TOPO_MLOPART_BUSID(busFn1, 0);
+  const int64_t enc7 = NCCL_TOPO_MLOPART_BUSID(busFn1, 7);
+  EXPECT_NE(enc7, enc0);
+  EXPECT_EQ(enc0 & ~NCCL_TOPO_MLOPART_MASK, busFn1);
+  EXPECT_EQ(enc7 & ~NCCL_TOPO_MLOPART_MASK, busFn1);
+}
+
 // ncclTopoAddXGMI() is only built on HIP/AMD platforms, so guard these tests
 // with the same macro; the #else branch keeps the suite linkable elsewhere.
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
@@ -415,6 +437,66 @@ TEST_F(TopoTest, DevModel_IndirectXgmiPath_IsFourHops) {
   ncclTopoFree(built);
 }
 
+TEST_F(TopoTest, GetSystemFromXml_MloPartOnNonzeroPhysicalFunction) {
+  const uint64_t host = 0x77;
+  struct ncclXmlNode* cpu = addSystemCpu(host);
+  struct ncclXmlNode* pci = addGpuPci(cpu, "0000:03:00.1", "gfx942", 0, 0, /*mloPart=*/0);
+  struct ncclXmlNode* gpu0 = nullptr;
+  ASSERT_EQ(xmlGetSub(pci, "gpu", &gpu0), ncclSuccess);
+  struct ncclXmlNode* gpu1 = addGpuUnderPci(pci, "gfx942", 1, 1, /*mloPart=*/1);
+  addGpuLink(gpu0, "0000:03:00.1", 1, PCI_ACCELERATOR_CLASS);
+  addGpuLink(gpu1, "0000:03:00.0", 1, PCI_ACCELERATOR_CLASS);
+
+  struct ncclTopoSystem* built = nullptr;
+  ASSERT_EQ(ncclTopoGetSystemFromXml(xml, &built, host), ncclSuccess);
+  ASSERT_NE(built, nullptr);
+  ASSERT_EQ(built->nodes[DEV].count, 2);
+  EXPECT_NE(findLink(built->nodes[DEV].nodes, built->nodes[DEV].nodes + 1), nullptr);
+  EXPECT_NE(findLink(built->nodes[DEV].nodes + 1, built->nodes[DEV].nodes), nullptr);
+  ncclTopoFree(built);
+}
+
+// CPX: eight HIP partitions (mlopart 0-7) hang off the physical function-0 PCI
+// node. Overlay bits must stay above the PCI function nibble.
+TEST_F(TopoTest, GetSystemFromXml_CpxEightMlopartsUnderPhysicalPci) {
+  const uint64_t host = 0xc0;
+  struct ncclXmlNode* cpu = addSystemCpu(host);
+  struct ncclXmlNode* pci = addGpuPci(cpu, "0000:0c:00.0", "gfx942", 0, 0, /*mloPart=*/0);
+  struct ncclXmlNode* gpu0 = nullptr;
+  ASSERT_EQ(xmlGetSub(pci, "gpu", &gpu0), ncclSuccess);
+  ASSERT_NE(gpu0, nullptr);
+  struct ncclXmlNode* gpus[NCCL_TOPO_MLOPART_DEV_MAX] = {};
+  gpus[0] = gpu0;
+  for (int p = 1; p < NCCL_TOPO_MLOPART_DEV_MAX; p++) {
+    gpus[p] = addGpuUnderPci(pci, "gfx942", p, p, p);
+  }
+  for (int p = 0; p < NCCL_TOPO_MLOPART_DEV_MAX; p++) {
+    char tgt[32];
+    snprintf(tgt, sizeof(tgt), "0000:0c:00.%d", (p + 1) % NCCL_TOPO_MLOPART_DEV_MAX);
+    addGpuLink(gpus[p], tgt, 1, PCI_ACCELERATOR_CLASS);
+  }
+
+  struct ncclTopoSystem* built = nullptr;
+  ASSERT_EQ(ncclTopoGetSystemFromXml(xml, &built, host), ncclSuccess);
+  ASSERT_NE(built, nullptr);
+  ASSERT_EQ(built->nodes[DEV].count, NCCL_TOPO_MLOPART_DEV_MAX);
+  ASSERT_EQ(built->nodes[GPU].count, NCCL_TOPO_MLOPART_DEV_MAX);
+
+  int64_t bus0 = 0;
+  ASSERT_EQ(busIdToInt64("0000:0c:00.0", &bus0), ncclSuccess);
+  for (int p = 0; p < NCCL_TOPO_MLOPART_DEV_MAX; p++) {
+    struct ncclTopoNode* dev = nullptr;
+    ASSERT_EQ(ncclTopoGetNode(built, &dev, DEV,
+                              NCCL_TOPO_ID(0, NCCL_TOPO_MLOPART_BUSID(bus0, p))),
+              ncclSuccess);
+    ASSERT_NE(dev, nullptr);
+    EXPECT_EQ(NCCL_TOPO_ID_LOCAL_ID(dev->id) & 0xf, 0);
+    EXPECT_NE(findLink(dev, built->nodes[DEV].nodes + (p + 1) % NCCL_TOPO_MLOPART_DEV_MAX), nullptr);
+  }
+
+  ncclTopoFree(built);
+}
+
 #else // !(__HIP_PLATFORM_AMD__ || __HIPCC__)
 
 // ncclTopoAddXGMI() is not built on non-HIP platforms, so register one skipped
diff --git a/projects/rccl/test/graph/XmlTests.cpp b/projects/rccl/test/graph/XmlTests.cpp
index 49b4b663265d..99cd76ba977d 100644
--- a/projects/rccl/test/graph/XmlTests.cpp
+++ b/projects/rccl/test/graph/XmlTests.cpp
@@ -1750,4 +1750,4 @@ TEST_F(XmlTest, TestXmlWarnings_EdgeCaseSuccess) {
   EXPECT_NE(systemNode, nullptr);
 
   free(xml);
-}
\ No newline at end of file
+}
diff --git a/projects/rccl/test/host/fakes/init_fakes.cc b/projects/rccl/test/host/fakes/init_fakes.cc
index d225c6a7324f..c94b489f8d9c 100644
--- a/projects/rccl/test/host/fakes/init_fakes.cc
+++ b/projects/rccl/test/host/fakes/init_fakes.cc
@@ -177,6 +177,12 @@ ncclResult_t ncclGpuGdrSupport(struct ncclComm*, int* gdrSupport) {
   if (gdrSupport) *gdrSupport = g_gdrSupportValue;
   return ncclSuccess;
 }
+// fillInfo MLOPart PCI-function fallback (init.cc ~1093): empty class keeps
+// isGpu=0 so existing tests' default busId=0 path stays a no-op.
+ncclResult_t ncclOsGetPciDeviceClassByBusId(const char* /*busId*/, char* deviceClass, size_t maxLen) {
+  if (deviceClass && maxLen > 0) deviceClass[0] = '\0';
+  return ncclSuccess;
+}
 ncclResult_t rocmLibraryInit(void) { return ncclSuccess; }
 uint64_t ncclOsGetPid() { return 4321; }
 // DMA-BUF export function pointer (dmaBufSupported gate): NULL -> unsupported.
diff --git a/projects/rccl/tools/topo_expl/src/topo_expl_impl.cpp b/projects/rccl/tools/topo_expl/src/topo_expl_impl.cpp
index 1a288ee7c5cc..d823ae088025 100644
--- a/projects/rccl/tools/topo_expl/src/topo_expl_impl.cpp
+++ b/projects/rccl/tools/topo_expl/src/topo_expl_impl.cpp
@@ -933,7 +933,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGatherInfo *a
   // AllGather3 - begin
   //NCCLCHECKGOTO(ncclCalloc(&allGather3Data, nranks), ret, fail);
   int idx;
-  NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, comm->busId, &idx));
+  NCCLCHECK(ncclTopoRankToIndex(comm->topo, rank, &idx, /*showWarn=*/true));
   allGather3Data[rank].nc = 2;
   if (comm->topo->nodes[GPU].count == comm->topo->nRanks &&
       IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx906") && allXgmi)