AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Allow intranode use of network connection (#383)

Bladeren bron 
* Allow intranode use of network connection

* Checking for graph for null pointer
This commit is contained in:
Wenkai Du
2021-06-08 07:37:59 -07:00
gecommit door GitHub
bovenliggende 820a53287f
commit a3a8c2d56b
7 gewijzigde bestanden met toevoegingen van 120 en 49 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
src/graph/rome_models.cc
+43 -35
Bestand weergeven
@@ -27,20 +27,16 @@ THE SOFTWARE.
#include <sys/time.h>
#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<ngpus; r++) {
int g = gpus[r];
@@ -380,10 +385,12 @@ ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct n
return ncclInternalError;
}
if (net) {
if (nets[0] >= 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);
src/graph/search.cc
+37 -2
Bestand weergeven
@@ -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; c<graph->nChannels; 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; i<ngpus; i++) {
int n = graph->intraNets[(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;
}
src/include/graph.h
+3
Bestand weergeven
@@ -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);
src/transport.cc
+7
Bestand weergeven
@@ -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<NTRANSPORTS; t++) {
if (connIndex == NCCL_CONN_IDX_P2P_NET && (t == TRANSPORT_SHM || (!xgmi && t == TRANSPORT_P2P)))
continue;
if (n1 >= 0 && n2 >= 0 && t != TRANSPORT_NET) continue;
struct ncclTransport *transport = ncclTransports+t;
struct ncclTransportComm* transportComm = type == 1 ? &transport->send : &transport->recv;
int ret = 0;
src/transport/net.cc
+12 -6
Bestand weergeven
@@ -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));
tools/topo_expl/model.cpp
+12 -6
Bestand weergeven
@@ -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,
tools/topo_expl/utils.cpp
+6
Bestand weergeven
@@ -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<NTRANSPORTS; t++) {
if (connIndex == NCCL_CONN_IDX_P2P_NET && (t == TRANSPORT_SHM || (!xgmi && t == TRANSPORT_P2P)))
continue;
if (n1 >= 0 && n2 >= 0 && t != TRANSPORT_NET) continue;
struct ncclTransport *transport = ncclTransports+t;
struct ncclTransportComm* transportComm = type == 1 ? &transport->send : &transport->recv;
int ret = 0;