From aa6d450f3506f712d176de59b6e665024452cca6 Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Mon, 22 Aug 2022 19:09:22 +0000 Subject: [PATCH] fix channelcount for multi-rank scenario [ROCm/rccl commit: 4141ec1151a5c5a1e6fb752d87c418b073e88130] --- projects/rccl/src/init.cc | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 670b257bc9..8238acab41 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -850,13 +850,19 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm int idx; NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, comm->busId, &idx)); allGather3Data[rank].nc = 2; - if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 906 && allXgmi) + if ( ((comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->virtualId == -1) || + (comm->topo->nodes[GPU].count <= comm->topo->nRanks && comm->virtualId != -1)) && + comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 906 && allXgmi) allGather3Data[rank].nc = 4; if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 908) allGather3Data[rank].nc = std::max(4/ringGraph.nChannels, 2); - if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G)) + if ( ((comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->virtualId == -1) || + (comm->topo->nodes[GPU].count <= comm->topo->nRanks && comm->virtualId != -1)) && + (comm->topo->type & RCCL_TOPO_CR8G)) allGather3Data[rank].nc = 4; - if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910) + if (((comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->virtualId == -1) || + (comm->topo->nodes[GPU].count <= comm->topo->nRanks && comm->virtualId != -1)) && + comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910) allGather3Data[rank].nc = 4; if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910) allGather3Data[rank].nc = std::max(allGather3Data[rank].nc, 4/ringGraph.nChannels);