diff --git a/src/graph/rome_models.cc b/src/graph/rome_models.cc index 29a0df3207..37a324b638 100755 --- a/src/graph/rome_models.cc +++ b/src/graph/rome_models.cc @@ -27,20 +27,16 @@ THE SOFTWARE. #include #include "rome_models.h" -#define MAX_ROME_CPUS 8 -#define MAX_ROME_GPUS 16 -#define MAX_ROME_NICS 16 - struct rcclRomeModel { int nGpus; int nCpus; int nNics; int nLinks; - int64_t gpuIds[MAX_ROME_GPUS]; - int64_t nicIds[MAX_ROME_NICS]; - int64_t gpuNuma[MAX_ROME_GPUS]; - int64_t nicNuma[MAX_ROME_NICS]; - int connMatrix[MAX_ROME_GPUS*MAX_ROME_GPUS]; + int64_t gpuIds[NCCL_TOPO_MAX_NODES]; + int64_t nicIds[NCCL_TOPO_MAX_NODES]; + int64_t gpuNuma[NCCL_TOPO_MAX_NODES]; + int64_t nicNuma[NCCL_TOPO_MAX_NODES]; + uint8_t connMatrix[NCCL_TOPO_MAX_NODES*NCCL_TOPO_MAX_NODES]; const char *pattern; const char *ringBase; }; @@ -319,13 +315,13 @@ static struct rcclRomeModel romeTopoModels[] = { * rings for multiple cases. */ ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int* gpu_map) { - int gpus[MAX_ROME_GPUS]; + int gpus[NCCL_TOPO_MAX_NODES]; int nChannels = 0; int gpu = 0; int offset = 0; - int status = 0; // 0 : between numbers, 1 : inside number, 2: start NET - int nets[2]; - int net = 0; + int status = 0; // 0 : between numbers, 1 : inside number, 2: start NET, 3: inside NET + int nets[NCCL_TOPO_MAX_NODES*2]; + int net_offset = 0, net_count = 0; int ngpus = system->nodes[GPU].count; int nnets = system->nodes[NET].count; do { @@ -336,29 +332,38 @@ ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct n } else { int digit = str[offset] - '0'; if (digit >= 0 && digit <= 9) { - if (status == 0) { - gpus[gpu] = digit; - status = 1; - } else if (status == 2) { - nets[net] = digit; - } - else{ - gpus[gpu] = gpus[gpu]*10+digit; + switch (status) { + case 0: + gpus[gpu] = digit; + status = 1; + break; + case 1: + gpus[gpu] = gpus[gpu]*10+digit; + break; + case 2: + nets[net_offset] = digit+'N'; + status = 3; + break; + case 3: + nets[net_offset] = (nets[net_offset]-'N')*10+digit+'N'; + break; } } else { if (status == 1) { gpu++; - if (gpu > MAX_ROME_GPUS) goto end; - } else if (status == 2) { - net++; - if (net > 2) goto end; + net_offset = 2*gpu-1; + if (gpu > NCCL_TOPO_MAX_NODES) goto end; + } else if (status == 2 || status == 3) { + net_offset++; + net_count++; + if (net_offset > ngpus*2) goto end; } status = 0; if (str[offset] == '|' || str[offset] == '\0') { // Ignore if ngpus doesn't match if (gpu != ngpus) goto newchannel; - // Ignore if nnets are not 0 or 2 - if (net && net != 2) goto newchannel; + // Ignore if net_count is not 0 or odd number + if (net_count && net_count%2) goto newchannel; for (int r=0; r= nnets || nets[1] >= nnets) goto newchannel; - graph->inter[nChannels*2] = system->nodes[NET].nodes[nets[0]].id; - graph->inter[nChannels*2+1] = system->nodes[NET].nodes[nets[1]].id; + if (net_count) { + memcpy(&graph->intraNets[ngpus*nChannels*2], nets, ngpus*2*sizeof(int)); + graph->nIntraChannels++; + if (nets[0]-'N' >= nnets || nets[ngpus*2-1]-'N' >= nnets) goto newchannel; + graph->inter[nChannels*2] = nets[0]-'N'; + graph->inter[nChannels*2+1] = nets[ngpus*2-1]-'N'; } else if (nnets) { graph->inter[nChannels*2] = system->nodes[NET].nodes[nChannels%nnets].id; graph->inter[nChannels*2+1] = system->nodes[NET].nodes[(nChannels+1)%nnets].id; @@ -391,7 +398,8 @@ ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct n nChannels++; newchannel: gpu = 0; - net = 0; + net_offset = 0; + net_count = 0; } } } @@ -514,14 +522,14 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo romeTopo->nNics = system->nodes[NET].count; romeTopo->nLinks = 0; // sort GPU devices by HIP device ID - struct ncclGpuIdHIP scores[MAX_ROME_GPUS]; + struct ncclGpuIdHIP scores[NCCL_TOPO_MAX_NODES]; for (int i = 0; i < romeTopo->nGpus; i ++) { scores[i].g = i; scores[i].dev = system->nodes[GPU].nodes[i].gpu.dev; } qsort(scores, romeTopo->nGpus, sizeof(struct ncclGpuIdHIP), cmpIds); // sort CPU devices by NUMA id - struct ncclCpuNuma cpu_scores[MAX_ROME_CPUS]; + struct ncclCpuNuma cpu_scores[NCCL_TOPO_MAX_NODES]; for (int i = 0; i < romeTopo->nCpus; i ++) { cpu_scores[i].c = i; cpu_scores[i].numa = system->nodes[CPU].nodes[i].id; @@ -683,7 +691,7 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph* // recognize system as Rome 4P2H even if no matching model if (ngpus > 4 && romeTopo.nLinks) system->type |= RCCL_TOPO_4P2H_ROME; - int g[MAX_ROME_GPUS]; + int g[NCCL_TOPO_MAX_NODES]; int time = 0; struct timeval tvs, tve; gettimeofday(&tvs, NULL); diff --git a/src/graph/search.cc b/src/graph/search.cc index 98edf7f105..8c839097aa 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -776,6 +776,8 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph graph->typeInter = PATH_PIX; graph->nChannels = 0; graph->sameChannels = 1; + graph->nIntraChannels = 0; + memset(graph->intraNets, 0, MAXCHANNELS*NCCL_TOPO_MAX_NODES*2*sizeof(int)); char* str = getenv("NCCL_GRAPH_FILE"); if (str) { @@ -944,6 +946,7 @@ done: int dupChannels = std::min(graph->nChannels*2, graph->maxChannels); memcpy(graph->intra+graph->nChannels*ngpus, graph->intra, (dupChannels-graph->nChannels)*ngpus*sizeof(int)); memcpy(graph->inter+graph->nChannels*2,graph->inter, (dupChannels-graph->nChannels)*2*sizeof(int)); + memcpy(graph->intraNets+graph->nChannels*ngpus*2, graph->intraNets, (dupChannels-graph->nChannels)*2*ngpus*sizeof(int)); graph->speedIntra /= DIVUP(dupChannels, graph->nChannels); graph->speedInter /= DIVUP(dupChannels, graph->nChannels); graph->nChannels = dupChannels; @@ -959,15 +962,25 @@ ncclResult_t ncclTopoPrintGraph(struct ncclTopoSystem* system, struct ncclTopoGr for (int c=0; cnChannels; c++) { sprintf(line, "%2d :", c); int offset = strlen(line); - if (system->nodes[NET].count > 0 && system->nodes[GPU].count != system->nRanks) { + if (system->nodes[NET].count > 0 && system->nodes[GPU].count != system->nRanks && !graph->nIntraChannels) { sprintf(line+offset, " %s/%d", topoNodeTypeStr[NET], graph->inter[2*c]); offset = strlen(line); } for (int i=0; iintraNets[(ngpus*c+i)*2]-'N'; + if(n >= 0 && n < system->nodes[NET].count) { + sprintf(line+offset, " NET/%d", n); + offset = strlen(line); + } sprintf(line+offset, " %s/%d", topoNodeTypeStr[GPU], graph->intra[ngpus*c+i]); offset = strlen(line); + n = graph->intraNets[(ngpus*c+i)*2+1]-'N'; + if(n >= 0 && n < system->nodes[NET].count) { + sprintf(line+offset, " NET/%d", n); + offset = strlen(line); + } } - if (system->nodes[NET].count > 0 && system->nodes[GPU].count != system->nRanks) { + if (system->nodes[NET].count > 0 && system->nodes[GPU].count != system->nRanks && !graph->nIntraChannels) { sprintf(line+offset, " %s/%d", topoNodeTypeStr[NET], graph->inter[2*c+1]); offset = strlen(line); } @@ -1003,3 +1016,25 @@ ncclResult_t ncclTopoGetNetDev(struct ncclTopoSystem* system, int rank, struct n } return ncclSuccess; } + +extern int64_t rcclParamP2pNetDisable(); + +ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, struct ncclTopoGraph* graph, int channelId, int type, int* dev) { + *dev = -1; + if (graph && graph->nIntraChannels && rcclParamP2pNetDisable() == 0) { + int n1 = -1; + int ngpus = system->nodes[GPU].count; + int nnets = system->nodes[NET].count; + int chan = channelId%graph->nIntraChannels; + for (int i = 0; i < ngpus; i++) { + if (graph->intra[ngpus*chan+i] == rank) { + n1 = graph->intraNets[(ngpus*chan+i)*2+type]-'N'; + break; + } + } + if (n1 >= 0 && n1 < nnets) { + *dev = n1; + } + } + return ncclSuccess; +} diff --git a/src/include/graph.h b/src/include/graph.h index 0e0929d278..4fa452e7d2 100644 --- a/src/include/graph.h +++ b/src/include/graph.h @@ -33,6 +33,7 @@ ncclResult_t ncclTopoGetNvbGpus(struct ncclTopoSystem* system, int rank, int* nr ncclResult_t ncclTopoGetNetDev(struct ncclTopoSystem* system, int rank, struct ncclTopoGraph* graph, int channelId, int rr, int* net); ncclResult_t ncclTopoCheckP2p(struct ncclTopoSystem* system, int64_t id1, int64_t id2, int* p2p, int *read, int* intermediateRank); ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* topo, int64_t busId, int netDev, int read, int* useGdr); +ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, struct ncclTopoGraph* graph, int channelId, int type, int* dev); // Set CPU affinity ncclResult_t ncclTopoSetAffinity(struct ncclTopoSystem* system, int rank); @@ -78,6 +79,8 @@ struct ncclTopoGraph { int nHops; int intra[MAXCHANNELS*NCCL_TOPO_MAX_NODES]; int inter[MAXCHANNELS*2]; + int nIntraChannels; + int intraNets[MAXCHANNELS*NCCL_TOPO_MAX_NODES*2]; }; ncclResult_t ncclTopoCompute(struct ncclTopoSystem* system, struct ncclTopoGraph* graph); diff --git a/src/transport.cc b/src/transport.cc index 5207212b71..76e76d1c79 100644 --- a/src/transport.cc +++ b/src/transport.cc @@ -36,11 +36,18 @@ static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* struct ncclPeerInfo* peerInfo = comm->peerInfo+peer; struct ncclConnector* connector = (type == 1) ? comm->channels[channelId].peers[peer].send + connIndex : comm->channels[channelId].peers[peer].recv + connIndex; + + // handle intra-node network connections + int n1, n2; + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, comm->rank, graph, channelId, (type == 1) ? 1 : 0, &n1)); + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, peer, graph, channelId, (type == 1) ? 0 : 1, &n2)); + int xgmi; NCCLCHECK(connectedByXGMI(&xgmi, comm->topo, myInfo, peerInfo)); for (int t=0; t= 0 && n2 >= 0 && t != TRANSPORT_NET) continue; struct ncclTransport *transport = ncclTransports+t; struct ncclTransportComm* transportComm = type == 1 ? &transport->send : &transport->recv; int ret = 0; diff --git a/src/transport/net.cc b/src/transport/net.cc index ce7aaedc22..204cf8e645 100644 --- a/src/transport/net.cc +++ b/src/transport/net.cc @@ -78,9 +78,12 @@ ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st send->conn.shared = resources->shared = ncclParamNetSharedBuffers() != -2 ? ncclParamNetSharedBuffers() : graph ? 0 : 1; send->proxyAppendPtr = send->conn.shared ? comm->proxyState.sharedBuffs.proxyAppend+2*channelId+1 : &send->proxyAppend; - // Send/Receive: Round-robin NICs based on the receiver's CUDA device - int nicRR = comm->peerInfo[peerInfo->rank].cudaDev; - NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &resources->netDev)); + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, myInfo->rank, graph, channelId, 1, &resources->netDev)); + if (resources->netDev < 0) { + // Send/Receive: Round-robin NICs based on the receiver's CUDA device + int nicRR = comm->peerInfo[peerInfo->rank].cudaDev; + NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &resources->netDev)); + } NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, resources->netDev, 1, &resources->useGdr)); NCCLCHECK(ncclCudaHostCalloc(&resources->sendMem, 1)); @@ -143,9 +146,12 @@ ncclResult_t netRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st recv->conn.shared = resources->shared = ncclParamNetSharedBuffers() != -2 ? ncclParamNetSharedBuffers() : graph ? 0 : 1; recv->proxyAppendPtr = recv->conn.shared ? comm->proxyState.sharedBuffs.proxyAppend+2*channelId : &recv->proxyAppend; - // Send/Receive: Round-robin NICs based on the receiver's CUDA device - int nicRR = comm->cudaDev; - NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &resources->netDev)); + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, myInfo->rank, graph, channelId, 0, &resources->netDev)); + if (resources->netDev < 0) { + // Send/Receive: Round-robin NICs based on the receiver's CUDA device + int nicRR = comm->cudaDev; + NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &resources->netDev)); + } NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, resources->netDev, 0, &resources->useGdr)); NCCLCHECK(ncclCudaHostCalloc(&resources->sendMem, 1)); diff --git a/tools/topo_expl/model.cpp b/tools/topo_expl/model.cpp index 58a0cd0d06..8b6874689c 100644 --- a/tools/topo_expl/model.cpp +++ b/tools/topo_expl/model.cpp @@ -151,9 +151,12 @@ ncclResult_t netCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* send, int channelId, int connIndex) { int netDev, useGdr = 0; - // Send/Receive: Round-robin NICs based on the receiver's CUDA device - int nicRR = comm->peerInfo[peerInfo->rank].cudaDev; - NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &netDev)); + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, myInfo->rank, graph, channelId, 1, &netDev)); + if (netDev < 0) { + // Send/Receive: Round-robin NICs based on the receiver's CUDA device + int nicRR = comm->peerInfo[peerInfo->rank].cudaDev; + NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &netDev)); + } NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netDev, 1, &useGdr)); INFO(NCCL_INIT|NCCL_NET,"Ring %02d : %d[%lx] -> %d[%lx] [send] via NET/%s/%d%s", channelId, myInfo->rank, myInfo->busId, peerInfo->rank, peerInfo->busId, ncclNetName(), netDev, @@ -166,9 +169,12 @@ NCCL_PARAM(NetGdrLevel, "NET_GDR_LEVEL", PATH_PHB); ncclResult_t netRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* recv, int channelId, int connIndex) { int netDev, useGdr = 0; - // Send/Receive: Round-robin NICs based on the receiver's CUDA device - int nicRR = comm->cudaDev; - NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &netDev)); + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, myInfo->rank, graph, channelId, 0, &netDev)); + if (netDev < 0) { + // Send/Receive: Round-robin NICs based on the receiver's CUDA device + int nicRR = comm->cudaDev; + NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, nicRR, &netDev)); + } NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netDev, 0, &useGdr)); INFO(NCCL_INIT|NCCL_NET,"Ring %02d : %d[%lx] -> %d[%lx] [receive] via NET/%s/%d%s", channelId, peerInfo->rank, peerInfo->busId, myInfo->rank, myInfo->busId, ncclNetName(), netDev, diff --git a/tools/topo_expl/utils.cpp b/tools/topo_expl/utils.cpp index e4ba62c19b..1f0fe72fb1 100644 --- a/tools/topo_expl/utils.cpp +++ b/tools/topo_expl/utils.cpp @@ -538,11 +538,17 @@ static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* struct ncclPeerInfo* peerInfo = comm->peerInfo+peer; struct ncclConnector* connector = (type == 1) ? comm->channels[channelId].peers[peer].send + connIndex : comm->channels[channelId].peers[peer].recv + connIndex; + // handle intra-node network connections + int n1, n2; + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, comm->rank, graph, channelId, (type == 1) ? 1 : 0, &n1)); + NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, peer, graph, channelId, (type == 1) ? 0 : 1, &n2)); + int xgmi; NCCLCHECK(connectedByXGMI(&xgmi, comm->topo, myInfo, peerInfo)); for (int t=0; t= 0 && n2 >= 0 && t != TRANSPORT_NET) continue; struct ncclTransport *transport = ncclTransports+t; struct ncclTransportComm* transportComm = type == 1 ? &transport->send : &transport->recv; int ret = 0;