diff --git a/src/graph/connect.cc b/src/graph/connect.cc
index 96d222ad16..6c681c397b 100644
--- a/src/graph/connect.cc
+++ b/src/graph/connect.cc
@@ -626,7 +626,8 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa
 
   // Only use full MAXCHANNELS for gfx94x
   int maxChannels = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ?
-    (comm->topo->nodes[GPU].nodes[0].gpu.cu == 80 ? 80 : MAXCHANNELS) : 2*CHANNEL_LIMIT;
+    ((comm->topo->nodes[GPU].nodes[0].gpu.cu == 80 || comm->topo->nodes[GPU].nodes[0].gpu.cu == 20)
+      ? comm->topo->nodes[GPU].nodes[0].gpu.cu : MAXCHANNELS) : 2*CHANNEL_LIMIT;
 
   if (graphs[NCCL_ALGO_RING]->nIntraChannels > 0 || comm->nNodes > 1) {
     maxChannels = std::min(64, maxChannels);
diff --git a/src/graph/xml.cc b/src/graph/xml.cc
index 3cfc7437b9..fc594380c4 100644
--- a/src/graph/xml.cc
+++ b/src/graph/xml.cc
@@ -347,7 +347,9 @@ static void memcpylower(char* dst, const char* src, const size_t size) {
 static ncclResult_t getPciPath(const char* busId, char** path) {
   char busPath[] = "/sys/class/pci_bus/0000:00/../../0000:00:00.0";
   memcpylower(busPath+sizeof("/sys/class/pci_bus/")-1, busId, BUSID_REDUCED_SIZE-1);
-  memcpylower(busPath+sizeof("/sys/class/pci_bus/0000:00/../../")-1, busId, BUSID_SIZE-1);
+  memcpylower(busPath+sizeof("/sys/class/pci_bus/0000:00/../../")-1, busId, BUSID_SIZE-2);
+  // override PCIe device function ID in CPX mode
+  busPath[sizeof("/sys/class/pci_bus/0000:00/../../")+BUSID_SIZE-3] = '0';
   *path = realpath(busPath, NULL);
   if (*path == NULL) {
     WARN("Could not find real path of %s", busPath);
diff --git a/src/graph/xml.h b/src/graph/xml.h
index 50773dc83a..95baba677c 100644
--- a/src/graph/xml.h
+++ b/src/graph/xml.h
@@ -18,7 +18,7 @@
 #define MAX_STR_LEN 255
 #define MAX_ATTR_COUNT 16
 #define MAX_SUBS 64	//Changed the value from 32 to 64 for CPX mode
-#define MAX_NODES 1024
+#define MAX_NODES 4096 //Changed the value from 1024 to 4096 for CPX mode
 
 #define NODE_TYPE_NONE 0
 #define NODE_TYPE_OPEN 1
diff --git a/src/misc/rocm_smi_wrap.cc b/src/misc/rocm_smi_wrap.cc
index 3589f41906..db40f87784 100644
--- a/src/misc/rocm_smi_wrap.cc
+++ b/src/misc/rocm_smi_wrap.cc
@@ -172,7 +172,8 @@ ncclResult_t rocm_smi_getLinkInfo(int srcIndex, int dstIndex, RSMI_IO_LINK_TYPE*
     if (rcclParamUseRocmSmiLib()) {
       ROCMSMICHECK(rsmi_topo_get_link_type(srcIndex, dstIndex, &rsmi_hops, rsmi_type));
       ROCMSMICHECK(rsmi_topo_get_link_weight(srcIndex, dstIndex, &rsmi_weight));
-      if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && rsmi_weight == 15) {
+      if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && (rsmi_weight == 15 ||
+        rsmi_weight == 41 || rsmi_weight == 13)) {
 	uint64_t min_bw = 0, max_bw = 0;
 	*hops = 1;
 #if defined USE_ROCM_SMI64CONFIG && rocm_smi_VERSION_MAJOR >= 5
@@ -189,7 +190,8 @@ ncclResult_t rocm_smi_getLinkInfo(int srcIndex, int dstIndex, RSMI_IO_LINK_TYPE*
       ARSMICHECK(ARSMI_topo_get_link_info(srcIndex, dstIndex, &tinfo));
 
       *rsmi_type  = (RSMI_IO_LINK_TYPE) tinfo.type;
-      if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && tinfo.weight == 15) {
+      if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && (tinfo.weight == 15 ||
+        tinfo.weight == 41 || tinfo.weight == 13)) {
 	*hops = 1;
 	if (tinfo.max_bandwidth && tinfo.min_bandwidth)
 	  *count = tinfo.max_bandwidth/tinfo.min_bandwidth;