@@ -74,8 +74,10 @@ struct ncclTransport {
|
||||
ncclResult_t ncclTransportP2pConnect(struct ncclComm* comm, int channelId, int nrecv, int* peerRecv, int nsend, int* peerSend, int connIndex);
|
||||
ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, int connIndex, int* highestTransportType=NULL);
|
||||
|
||||
#if CUDART_VERSION >= 12010
|
||||
ncclResult_t ncclNvlsSetup(struct ncclComm* comm);
|
||||
ncclResult_t ncclNvlsFree(struct ncclComm* comm);
|
||||
#endif
|
||||
|
||||
enum { collNetRecv=0, collNetSend=1 };
|
||||
int ncclTransportCollNetSetup(struct ncclComm* comm, struct ncclTopoGraph* collNetGraph, struct ncclChannel* channel, int masterRank, int masterPeer, int collNetGraphChannelId, int type);
|
||||
|
||||
@@ -373,7 +373,9 @@ static ncclResult_t commFree(ncclComm_t comm) {
|
||||
NCCLCHECK(ncclStrongStreamDestruct(&comm->deviceStream));
|
||||
}
|
||||
|
||||
#if CUDART_VERSION >= 12010
|
||||
if (comm->nvlsSupport) NCCLCHECK(ncclNvlsFree(comm));
|
||||
#endif
|
||||
|
||||
struct ncclDestructor* dtor = comm->destructorHead;
|
||||
while (dtor != nullptr) {
|
||||
@@ -1268,7 +1270,9 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
|
||||
// Check if we can setup CollNet
|
||||
if (comm->collNetSupport > 0) collNetTrySetup(comm, &collNetGraph);
|
||||
|
||||
//NCCLCHECKGOTO(ncclNvlsSetup(comm), ret, fail);
|
||||
#if CUDART_VERSION >= 12010
|
||||
NCCLCHECKGOTO(ncclNvlsSetup(comm), ret, fail);
|
||||
#endif
|
||||
|
||||
TRACE(NCCL_INIT, "rank %d nranks %d - CONNECTED %d RINGS AND TREES", rank, nranks, comm->nChannels);
|
||||
|
||||
|
||||
Ссылка в новой задаче
Block a user