diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc index db74694e1e..7c29920d1e 100644 --- a/projects/rccl/src/graph/paths.cc +++ b/projects/rccl/src/graph/paths.cc @@ -573,7 +573,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) { } } - if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_4P2H_ROME) && !(comm->topo->type & RCCL_TOPO_GDR_ALL)) { + if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_4P2H_ROME) && !(comm->topo->type & RCCL_TOPO_GDR_ALL) && (comm->topo->nodes[GPU].nodes[0].gpu.gcn != 910)) { // Adjust P2P channels on Rome comm->p2pnChannelsPerPeer = 2; comm->p2pnChannels = 2; diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index d45cd7f5e3..6b90197f79 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -949,6 +949,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm allGather3Data[rank].nc = comm->nChannels*4; if (comm->topo->nodes[GPU].count != comm->topo->nRanks && comm->topo->nodes[NET].count && (comm->topo->type & RCCL_TOPO_4P2H_ROME)) allGather3Data[rank].nc = (comm->topo->nodes[NET].count > 3 ? 2 : 4)*comm->topo->nodes[NET].count; + if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910) + allGather3Data[rank].nc = comm->nChannels*6; allGather3Data[rank].tree.pattern = treeGraph.pattern; allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels; allGather3Data[rank].tree.speedIntra = treeGraph.speedIntra;