From 2c0c02b211b0b9929243af8be79426edfe506472 Mon Sep 17 00:00:00 2001 From: Nilesh M Negi Date: Wed, 16 Jul 2025 09:07:02 -0500 Subject: [PATCH] [GRAPH] Match maxChannels for gfx942 CUs (#1302) [ROCm/rccl commit: 6632183efe9d283f4356422571dcc41cedd4ebe8] --- projects/rccl/src/graph/connect.cc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/rccl/src/graph/connect.cc b/projects/rccl/src/graph/connect.cc index 17a7ff85db..c26a775555 100644 --- a/projects/rccl/src/graph/connect.cc +++ b/projects/rccl/src/graph/connect.cc @@ -804,10 +804,10 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa } } - // Only use full MAXCHANNELS for gfx94x and gfx950 - maxChannels = (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) ? - ((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 == 38) - ? comm->topo->nodes[GPU].nodes[0].gpu.cu : MAXCHANNELS) : 2*CHANNEL_LIMIT; + // Only use full MAXCHANNELS for gfx942 (MI300X) and gfx950 + maxChannels = (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || + IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) + ? std::min(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);