AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

Fix topo explorer's nccl 2.23 compatibility (#1623)

Parcourir la source 
* Fix compiler issues due to broken compatibility 

* Fix segfault and pass rank instead of busid and add a pointer to cover a new algorithm
Cette révision appartient à :
Mustafa Abduljabbar
2025-04-02 09:47:29 -04:00
révisé par GitHub
Parent ffe255d285
révision aace4e27f8
2 fichiers modifiés avec 12 ajouts et 10 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
tools/topo_expl/model.cpp
+7 -7
Voir le fichier
@@ -78,7 +78,7 @@ int busIdToCudaDev(int64_t busId) {
static int useMemcpy = 0; static int useMemcpy = 0;
/* Determine if two peers can communicate with P2P */ /* Determine if two peers can communicate with P2P */
ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) { ncclResult_t p2pCanConnect(int* ret, struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
if (!info1->hasFineGrain || !info2->hasFineGrain) { if (!info1->hasFineGrain || !info2->hasFineGrain) {
*ret = 0; *ret = 0;
return ncclSuccess; return ncclSuccess;
@@ -92,7 +92,7 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop
// Check topology / p2p level. // Check topology / p2p level.
int intermediateRank; int intermediateRank;
NCCLCHECK(ncclTopoCheckP2p(topo, info1->busId, info2->busId, ret, NULL, &intermediateRank)); NCCLCHECK(ncclTopoCheckP2p(comm->topo, info1->rank, info2->rank, ret, NULL, &intermediateRank));
if (*ret == 0) return ncclSuccess; if (*ret == 0) return ncclSuccess;
if (intermediateRank != -1) { if (intermediateRank != -1) {
if (useMemcpy) *ret = 0; if (useMemcpy) *ret = 0;
@@ -101,7 +101,7 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop
// Check if NET would work better // Check if NET would work better
int useNet = 0; int useNet = 0;
NCCLCHECK(ncclTopoCheckNet(topo, info1->busId, info2->busId, &useNet)); NCCLCHECK(ncclTopoCheckNet(comm->topo, info1->rank, info2->rank, &useNet));
if (useNet) { if (useNet) {
*ret = 0; *ret = 0;
return ncclSuccess; return ncclSuccess;
@@ -146,12 +146,12 @@ struct ncclTransport p2pTransport = {
NCCL_PARAM(ShmDisable, "SHM_DISABLE", 0); NCCL_PARAM(ShmDisable, "SHM_DISABLE", 0);
/* Determine if two peers can communicate with SHM */ /* Determine if two peers can communicate with SHM */
ncclResult_t shmCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) { ncclResult_t shmCanConnect(int* ret, struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
*ret = 0; *ret = 0;
if (ncclParamShmDisable() == 1) return ncclSuccess; if (ncclParamShmDisable() == 1) return ncclSuccess;
int useNet = 0; int useNet = 0;
NCCLCHECK(ncclTopoCheckNet(topo, info1->busId, info2->busId, &useNet)); NCCLCHECK(ncclTopoCheckNet(comm->topo, info1->rank, info2->rank, &useNet));
if (useNet) return ncclSuccess; if (useNet) return ncclSuccess;
// Same host? // Same host?
@@ -197,7 +197,7 @@ struct setupReq {
}; };
/* Determine if two peers can communicate with NET */ /* Determine if two peers can communicate with NET */
ncclResult_t netCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) { ncclResult_t netCanConnect(int* ret, struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
*ret = 1; *ret = 1;
return ncclSuccess; return ncclSuccess;
} }
@@ -257,7 +257,7 @@ struct ncclTransport netTransport = {
}; };
/* Determine if two peers can communicate with NET */ /* Determine if two peers can communicate with NET */
ncclResult_t collNetCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) { ncclResult_t collNetCanConnect(int* ret, struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
*ret = 1; *ret = 1;
return ncclSuccess; return ncclSuccess;
} }
tools/topo_expl/utils.cpp
+5 -3
Voir le fichier
@@ -353,7 +353,7 @@ static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph*
struct ncclTransport *transport = ncclTransports[t]; struct ncclTransport *transport = ncclTransports[t];
struct ncclTransportComm* transportComm = type == 1 ? &transport->send : &transport->recv; struct ncclTransportComm* transportComm = type == 1 ? &transport->send : &transport->recv;
int ret = 0; int ret = 0;
NCCLCHECK(transport->canConnect(&ret, comm->topo, graph, myInfo, peerInfo)); NCCLCHECK(transport->canConnect(&ret, comm, graph, myInfo, peerInfo));
if (ret) { if (ret) {
connector->transportComm = transportComm; connector->transportComm = transportComm;
NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex)); NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex));
@@ -721,7 +721,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGatherInfo *a
//struct ncclTopoGraph treeGraph; //struct ncclTopoGraph treeGraph;
//struct ncclTopoGraph collNetGraph; //struct ncclTopoGraph collNetGraph;
//struct ncclTopoGraph nvlsGraph; //struct ncclTopoGraph nvlsGraph;
struct ncclTopoGraph* graphs[] = { &treeGraph, &ringGraph, &collNetGraph, &collNetGraph, &nvlsGraph, &nvlsGraph }; struct ncclTopoGraph* graphs[] = { &treeGraph, &ringGraph, &collNetGraph, &collNetGraph, &nvlsGraph, &nvlsGraph, &nvlsGraph};
int nChannelsOrig; int nChannelsOrig;
struct ncclTopoRanks** allTopoRanks = NULL; struct ncclTopoRanks** allTopoRanks = NULL;
@@ -963,6 +963,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGatherInfo *a
} }
comm->nChannels = std::min(treeGraph.nChannels, ringGraph.nChannels); comm->nChannels = std::min(treeGraph.nChannels, ringGraph.nChannels);
TRACE(NCCL_INIT,"treeGraph.nChannels: %d , ringGraph.nChannels: %d", treeGraph.nChannels, ringGraph.nChannels);
NCCLCHECKGOTO(ncclTopoPreset(comm, graphs, &allGather3Data[rank].topoRanks), ret, fail); NCCLCHECKGOTO(ncclTopoPreset(comm, graphs, &allGather3Data[rank].topoRanks), ret, fail);
fail: fail:
return ret; return ret;
@@ -975,7 +976,7 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGatherInfo *a
int nranks = comm->nRanks; int nranks = comm->nRanks;
cpu_set_t affinitySave; cpu_set_t affinitySave;
struct ncclTopoGraph* graphs[] = { &treeGraph, &ringGraph, &collNetGraph, &collNetGraph, &nvlsGraph, &nvlsGraph }; struct ncclTopoGraph* graphs[] = { &treeGraph, &ringGraph, &collNetGraph, &collNetGraph, &nvlsGraph, &nvlsGraph, &nvlsGraph};
int nChannelsOrig; int nChannelsOrig;
struct ncclTopoRanks** allTopoRanks = NULL; struct ncclTopoRanks** allTopoRanks = NULL;
@@ -1090,6 +1091,7 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGatherInfo *a
} }
NCCLCHECKGOTO(ncclCalloc(&rings, nranks*MAXCHANNELS), ret, fail); NCCLCHECKGOTO(ncclCalloc(&rings, nranks*MAXCHANNELS), ret, fail);
TRACE(NCCL_INIT, "rank %d nranks %d nchannels %d", rank, nranks, comm->nChannels);
NCCLCHECKGOTO(ncclTopoPostset(comm, nodesFirstRank, nodesTreePatterns, allTopoRanks, rings, graphs, NULL, nc), ret, fail); NCCLCHECKGOTO(ncclTopoPostset(comm, nodesFirstRank, nodesTreePatterns, allTopoRanks, rings, graphs, NULL, nc), ret, fail);
if (comm->topo->treeDefined) NCCLCHECK(ncclTreeBasePostset(comm, &treeGraph)); if (comm->topo->treeDefined) NCCLCHECK(ncclTreeBasePostset(comm, &treeGraph));