Merge pull request #1022 from PedramAlizadeh/sync_nccl_2.18.6
Sync to nccl 2.18.6
[ROCm/rccl commit: aa5c84c997]
Этот коммит содержится в:
Pedram Alizadeh
GitHub
@@ -1,6 +1,6 @@
|
||||
##### version
|
||||
NCCL_MAJOR := 2
|
||||
NCCL_MINOR := 18
|
||||
NCCL_PATCH := 3
|
||||
NCCL_PATCH := 6
|
||||
NCCL_SUFFIX :=
|
||||
PKG_REVISION := 1
|
||||
|
||||
@@ -42,9 +42,11 @@ ncclResult_t initChannel(struct ncclComm* comm, int channelId) {
|
||||
/* channel->devPeers is not shared, so just free it when calling commFree() */
|
||||
NCCLCHECK(ncclCudaCallocAsync(&channel->devPeers, nPeers, sharedRes->deviceStream.cudaStream));
|
||||
ncclCommPushCudaFree(comm, channel->devPeers);
|
||||
NCCLCHECK(ncclCalloc(&channel->devPeersHostPtr, nPeers));
|
||||
for (int r = 0; r < nRanks; r++) {
|
||||
uintptr_t addr = (uintptr_t)(comm->sharedRes->devPeers[channelId] + comm->topParentRanks[r]);
|
||||
NCCLCHECK(ncclCudaMemcpyAsync((uintptr_t*)(channel->devPeers + r), (uintptr_t*)&addr, 1, sharedRes->deviceStream.cudaStream));
|
||||
channel->devPeersHostPtr[r] = (struct ncclDevChannelPeer*)addr;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -52,9 +54,9 @@ ncclResult_t initChannel(struct ncclComm* comm, int channelId) {
|
||||
NCCLCHECK(ncclCudaCallocAsync(&channel->devRingUserRanks, nRanks, sharedRes->deviceStream.cudaStream));
|
||||
ncclCommPushCudaFree(comm, channel->devRingUserRanks);
|
||||
|
||||
/* guarantee addr has been copied into channel->devPeers */
|
||||
NCCLCHECK(ncclStrongStreamSynchronize(&sharedRes->deviceStream));
|
||||
NCCLCHECK(ncclStrongStreamRelease(ncclCudaGraphNone(), &sharedRes->deviceStream));
|
||||
CUDACHECK(hipEventRecord(sharedRes->deviceStream.scratchEvent, sharedRes->deviceStream.cudaStream));
|
||||
CUDACHECK(hipStreamWaitEvent(sharedRes->deviceStream.cudaStream, sharedRes->deviceStream.scratchEvent, 0));
|
||||
|
||||
return ncclSuccess;
|
||||
}
|
||||
@@ -79,6 +81,7 @@ ncclResult_t initNvlsChannel(struct ncclComm* comm, int channelId, struct ncclCo
|
||||
uintptr_t addr = (uintptr_t)(parent->channels[channelId].nvlsDevPeers + tr);
|
||||
channel->peers[comm->nRanks + 1 + r] = parent->channels[channelId].nvlsPeers + tr;
|
||||
NCCLCHECK(ncclCudaMemcpyAsync((uintptr_t*)(channel->devPeers + comm->nRanks + 1 + r), (uintptr_t*)&addr, 1, sharedRes->deviceStream.cudaStream));
|
||||
channel->devPeersHostPtr[comm->nRanks + 1 + r] = (struct ncclDevChannelPeer*)addr;
|
||||
ncclAtomicRefCountIncrement(&parent->channels[channelId].nvlsPeers[tr].refCount);
|
||||
}
|
||||
} else {
|
||||
@@ -88,10 +91,12 @@ ncclResult_t initNvlsChannel(struct ncclComm* comm, int channelId, struct ncclCo
|
||||
uintptr_t addr = (uintptr_t)(channel->nvlsDevPeers + r);
|
||||
channel->peers[comm->nRanks + 1 + r] = channel->nvlsPeers + r;
|
||||
NCCLCHECK(ncclCudaMemcpyAsync((uintptr_t*)(channel->devPeers + comm->nRanks + 1 + r), (uintptr_t*)&addr, 1, sharedRes->deviceStream.cudaStream));
|
||||
channel->devPeersHostPtr[comm->nRanks + 1 + r] = (struct ncclDevChannelPeer*)addr;
|
||||
ncclAtomicRefCountIncrement(&channel->nvlsPeers[r].refCount);
|
||||
}
|
||||
}
|
||||
|
||||
NCCLCHECK(ncclStrongStreamSynchronize(&sharedRes->deviceStream));
|
||||
NCCLCHECK(ncclStrongStreamRelease(ncclCudaGraphNone(), &sharedRes->deviceStream));
|
||||
|
||||
return ncclSuccess;
|
||||
@@ -116,6 +121,7 @@ ncclResult_t initCollnetChannel(struct ncclComm* comm, int channelId, struct ncc
|
||||
addr = (uintptr_t)parent->channels[channelId].collnetDevPeers;
|
||||
channel->peers[comm->nRanks] = parent->channels[channelId].collnetPeers;
|
||||
NCCLCHECK(ncclCudaMemcpyAsync((uintptr_t*)(channel->devPeers + comm->nRanks), (uintptr_t*)&addr, 1, sharedRes->deviceStream.cudaStream));
|
||||
channel->devPeersHostPtr[comm->nRanks] = (struct ncclDevChannelPeer*)addr;
|
||||
ncclAtomicRefCountIncrement(&parent->channels[channelId].collnetPeers->refCount);
|
||||
} else {
|
||||
NCCLCHECK(ncclCalloc(&channel->collnetPeers, 1));
|
||||
@@ -123,9 +129,11 @@ ncclResult_t initCollnetChannel(struct ncclComm* comm, int channelId, struct ncc
|
||||
addr = (uintptr_t)channel->collnetDevPeers;
|
||||
channel->peers[comm->nRanks] = channel->collnetPeers;
|
||||
NCCLCHECK(ncclCudaMemcpyAsync((uintptr_t*)(channel->devPeers + comm->nRanks), (uintptr_t*)&addr, 1, sharedRes->deviceStream.cudaStream));
|
||||
channel->devPeersHostPtr[comm->nRanks] = (struct ncclDevChannelPeer*)addr;
|
||||
ncclAtomicRefCountIncrement(&channel->collnetPeers->refCount);
|
||||
}
|
||||
|
||||
NCCLCHECK(ncclStrongStreamSynchronize(&sharedRes->deviceStream));
|
||||
NCCLCHECK(ncclStrongStreamRelease(ncclCudaGraphNone(), &sharedRes->deviceStream));
|
||||
|
||||
return ncclSuccess;
|
||||
@@ -158,5 +166,7 @@ ncclResult_t freeChannel(struct ncclChannel* channel, int nRanks, int collnetNRa
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
free(channel->devPeersHostPtr);
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
@@ -866,9 +866,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) {
|
||||
// fill the whole space of nChannels. To do so we mirror the bits in the
|
||||
// nChannels space.
|
||||
for (int c=0; c<comm->p2pnChannels; c++) {
|
||||
int mirror = 0;
|
||||
for (int b=1, mb=(comm->p2pnChannels>>1); b<comm->p2pnChannels; b<<=1, mb>>=1) if (c & b) mirror |= mb;
|
||||
comm->p2pChannels[c] = mirror;
|
||||
comm->p2pChannels[c] = mirrorBits(c, comm->p2pnChannels);
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
@@ -412,6 +412,28 @@ ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, in
|
||||
int* localNets;
|
||||
NCCLCHECK(ncclCalloc(&localNets, system->nodes[NET].count));
|
||||
|
||||
// First add the preferred NICs
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
if (gpu != -1 && gpu != g) continue;
|
||||
localNetCount = 0;
|
||||
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
|
||||
for (int c = 0;; c++) {
|
||||
int netId;
|
||||
NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &netId));
|
||||
NCCLCHECK(ncclTopoIdToIndex(system, NET, netId, localNets+localNetCount));
|
||||
if (localNetCount > 0 && localNets[localNetCount] == localNets[0]) break;
|
||||
localNetCount++;
|
||||
}
|
||||
// Append NICs to list
|
||||
for (int i=0; i<localNetCount; i++) {
|
||||
int n = localNets[i];
|
||||
int found = 0;
|
||||
while (nets[found] != n && found<netCount) found++;
|
||||
if (found == netCount) nets[netCount++] = n;
|
||||
}
|
||||
}
|
||||
|
||||
// Then add others satisfying typeInter
|
||||
for (int t=0; t <= typeInter; t++) {
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
if (gpu != -1 && gpu != g) continue;
|
||||
@@ -421,14 +443,6 @@ ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, in
|
||||
for (int n=0; n<system->nodes[NET].count; n++) {
|
||||
if (paths[n].type == t) localNets[localNetCount++] = n;
|
||||
}
|
||||
if (localNetCount == 0) continue;
|
||||
// Shuffle by gpu NVML device number so that GPUs on the same PCI switch
|
||||
// with multiple NICs don't use the same one as first choice.
|
||||
for (int r=0; r<system->nodes[GPU].nodes[g].gpu.dev % localNetCount; r++) {
|
||||
int net0 = localNets[0];
|
||||
for (int i=0; i<localNetCount-1; i++) localNets[i] = localNets[i+1];
|
||||
localNets[localNetCount-1] = net0;
|
||||
}
|
||||
// Append NICs to list
|
||||
for (int i=0; i<localNetCount; i++) {
|
||||
int n = localNets[i];
|
||||
@@ -566,7 +580,11 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
|
||||
// NVLS needs to balance on all NICs
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) {
|
||||
NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, -1, -1, nets[graph->nChannels]));
|
||||
if (graph->nChannels < netcount) {
|
||||
int gpu;
|
||||
NCCLCHECK(ncclTopoGetLocalGpu(system, system->nodes[NET].nodes[nets[graph->nChannels]].id, &gpu));
|
||||
if (gpu != -1) NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, -1, -1, gpu));
|
||||
}
|
||||
} else {
|
||||
if (graph->nChannels > 0) {
|
||||
// Try to replay the last channel
|
||||
|
||||
@@ -758,126 +758,61 @@ ncclResult_t ncclTopoGetSystem(struct ncclComm* comm, struct ncclTopoSystem** sy
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
static ncclResult_t getLocalNetMask(struct ncclTopoSystem* system, int g, uint64_t* localNetMask, int* type) {
|
||||
ncclResult_t ncclTopoGetLocal(struct ncclTopoSystem* system, int type, int index, int resultType, int** locals, int* localCount, int* pathType) {
|
||||
int minType = PATH_DIS;
|
||||
float maxBw = 0;
|
||||
int count = 0;
|
||||
int* nets;
|
||||
NCCLCHECK(ncclCalloc(&nets, system->nodes[NET].count));
|
||||
for (int n=0; n<system->nodes[NET].count; n++) {
|
||||
struct ncclTopoLinkList* path = system->nodes[NET].nodes[n].paths[GPU]+g;
|
||||
if (path->bw > maxBw || (path->bw == maxBw && path->type < minType)) {
|
||||
maxBw = path->bw;
|
||||
minType = path->type;
|
||||
if (type) *type = minType;
|
||||
NCCLCHECK(ncclCalloc(locals, system->nodes[resultType].count));
|
||||
struct ncclTopoLinkList* paths = system->nodes[type].nodes[index].paths[resultType];
|
||||
for (int i=0; i<system->nodes[resultType].count; i++) {
|
||||
if (paths[i].bw > maxBw || (paths[i].bw == maxBw && paths[i].type < minType)) {
|
||||
maxBw = paths[i].bw;
|
||||
minType = paths[i].type;
|
||||
if (pathType) *pathType = minType;
|
||||
count = 0;
|
||||
}
|
||||
if (path->bw == maxBw && path->type == minType) nets[count++] = system->nodes[NET].nodes[n].id;
|
||||
if (paths[i].bw == maxBw && paths[i].type == minType) (*locals)[count++] = i;
|
||||
}
|
||||
|
||||
*localNetMask = 0ULL;
|
||||
for (int n=0; n<count; n++) {
|
||||
if (nets[n] >= 64) return ncclInternalError;
|
||||
*localNetMask |= 1ULL<<nets[n];
|
||||
}
|
||||
free(nets);
|
||||
*localCount = count;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t ncclTopoGetLocalNet(struct ncclTopoSystem* system, int rank, int channelId, int* id) {
|
||||
uint64_t* localNetMasks;
|
||||
int ngpus = system->nodes[GPU].count;
|
||||
NCCLCHECK(ncclCalloc(&localNetMasks, ngpus));
|
||||
|
||||
// Fill localNetMasks for all GPUs.
|
||||
for (int g=0; g<ngpus; g++) {
|
||||
NCCLCHECK(getLocalNetMask(system, g, localNetMasks+g, NULL));
|
||||
}
|
||||
|
||||
// Find GPUs which have the same mask as rank, i.e. share the same local Nets.
|
||||
int gpu;
|
||||
NCCLCHECK(ncclTopoRankToIndex(system, rank, &gpu));
|
||||
int netLocalGpus = 0, netLocalGpu = 0;
|
||||
for (int g=0; g<ngpus; g++) {
|
||||
if (localNetMasks[g] == localNetMasks[gpu]) {
|
||||
if (g == gpu) netLocalGpu = netLocalGpus;
|
||||
netLocalGpus++;
|
||||
}
|
||||
}
|
||||
uint64_t localNetMask = localNetMasks[gpu];
|
||||
free(localNetMasks);
|
||||
if (localNetMask == 0) return ncclInternalError;
|
||||
|
||||
// Round robin on GPUs and channels
|
||||
int gIndex = 0, cId = 0, n = 0;
|
||||
while (1) {
|
||||
if (1ULL << n & localNetMask) {
|
||||
if (gIndex == netLocalGpu && cId == channelId) {
|
||||
*id = n;
|
||||
return ncclSuccess;
|
||||
}
|
||||
gIndex++;
|
||||
if (gIndex == netLocalGpus) {
|
||||
gIndex = 0;
|
||||
cId++;
|
||||
}
|
||||
}
|
||||
n = (n+1) % 64;
|
||||
int* localNets;
|
||||
int localNetCount;
|
||||
NCCLCHECK(ncclTopoGetLocal(system, GPU, gpu, NET, &localNets, &localNetCount, NULL));
|
||||
int* localGpus;
|
||||
int localGpuCount;
|
||||
NCCLCHECK(ncclTopoGetLocal(system, NET, localNets[0], GPU, &localGpus, &localGpuCount, NULL));
|
||||
int net = system->nodes[GPU].nodes[gpu].gpu.dev;
|
||||
if (isPow2(localNetCount)) net = mirrorBits(net, localNetCount);
|
||||
if (localNetCount == 0) {
|
||||
*id = -1;
|
||||
} else {
|
||||
net += channelId%(DIVUP(localNetCount,localGpuCount));
|
||||
*id = system->nodes[NET].nodes[localNets[net%localNetCount]].id;
|
||||
}
|
||||
free(localNets);
|
||||
free(localGpus);
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t ncclTopoGetLocalGpu(struct ncclTopoSystem* system, int net, int* gpuIndex) {
|
||||
int ngpus = system->nodes[GPU].count;
|
||||
int* gpus;
|
||||
NCCLCHECK(ncclCalloc(&gpus, ngpus));
|
||||
|
||||
// Find localNetMask which includes net with the most local GPUs.
|
||||
int netLocalGpus = 0, minType = PATH_DIS;
|
||||
uint64_t localNetMask = 0ULL;
|
||||
for (int g=0; g<ngpus; g++) {
|
||||
int type = PATH_DIS;
|
||||
uint64_t mask;
|
||||
NCCLCHECK(getLocalNetMask(system, g, &mask, &type));
|
||||
if ((1ULL<<net) & mask) {
|
||||
if (type < minType) {
|
||||
localNetMask = mask;
|
||||
netLocalGpus = 0;
|
||||
minType = type;
|
||||
}
|
||||
if (type == minType) {
|
||||
if (localNetMask && mask != localNetMask) {
|
||||
WARN("Gpus %d and %d both have a type of %d with net %d yet have different netMasks of %lx and %lx\n", g, gpus[netLocalGpus-1], minType, net, mask, localNetMask);
|
||||
free(gpus);
|
||||
return ncclInternalError;
|
||||
}
|
||||
gpus[netLocalGpus] = g;
|
||||
netLocalGpus++;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (localNetMask == 0ULL) {
|
||||
*gpuIndex = -1;
|
||||
free(gpus);
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Round robin on GPUs and channels
|
||||
int gIndex = 0, cId = 0, n = 0;
|
||||
while (1) {
|
||||
if (1ULL << n & localNetMask) {
|
||||
if (n == net) {
|
||||
*gpuIndex = gpus[gIndex];
|
||||
free(gpus);
|
||||
for (int c=0; c<MAXCHANNELS; c++) {
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
|
||||
int id;
|
||||
NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &id));
|
||||
if (net == id) {
|
||||
*gpuIndex = g;
|
||||
return ncclSuccess;
|
||||
}
|
||||
gIndex++;
|
||||
if (gIndex == netLocalGpus) {
|
||||
gIndex = 0;
|
||||
cId++;
|
||||
}
|
||||
}
|
||||
n = (n+1) % 64;
|
||||
}
|
||||
*gpuIndex = -1;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
/****************************/
|
||||
|
||||
@@ -241,4 +241,13 @@ static float ncclTopoXGMISpeed(const char* gcn) {
|
||||
#define ncclGetKernelIndex(p_comm) (0)
|
||||
#endif
|
||||
|
||||
// Mirror bits
|
||||
static bool isPow2(int val) {
|
||||
return (val & (val-1)) == 0;
|
||||
}
|
||||
static int mirrorBits(int val, int pow2) {
|
||||
int mirror = 0;
|
||||
for (int b=1, mb=(pow2>>1); b<pow2; b<<=1, mb>>=1) if (val & b) mirror |= mb;
|
||||
return mirror;
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -592,7 +592,18 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml*
|
||||
}
|
||||
}
|
||||
pciNode->parent = parent;
|
||||
parent->subs[parent->nSubs++] = pciNode;
|
||||
// Keep PCI sub devices ordered by PCI Bus ID (Issue #820)
|
||||
int subIndex = parent->nSubs;
|
||||
const char* newBusId;
|
||||
NCCLCHECK(xmlGetAttrStr(pciNode, "busid", &newBusId));
|
||||
for (int s=0; s<parent->nSubs; s++) {
|
||||
const char* busId;
|
||||
NCCLCHECK(xmlGetAttrStr(parent->subs[s], "busid", &busId));
|
||||
if (strcmp(newBusId, busId) < 0) { subIndex = s; break; }
|
||||
}
|
||||
for (int s = parent->nSubs; s > subIndex; s--) parent->subs[s] = parent->subs[s-1];
|
||||
parent->subs[subIndex] = pciNode;
|
||||
parent->nSubs++;
|
||||
}
|
||||
if (strcmp(parent->name, "pci") == 0) {
|
||||
NCCLCHECK(ncclTopoGetXmlFromSys(parent, xml));
|
||||
|
||||
@@ -129,6 +129,8 @@ struct ncclSharedResources {
|
||||
struct ncclChannel {
|
||||
struct ncclChannelPeer** peers;
|
||||
struct ncclDevChannelPeer** devPeers;
|
||||
/* devPeer pointer array used for host side access */
|
||||
struct ncclDevChannelPeer** devPeersHostPtr;
|
||||
struct ncclRing ring;
|
||||
int* devRingUserRanks;
|
||||
struct ncclTree tree;
|
||||
|
||||
@@ -672,7 +672,7 @@ static ncclResult_t devCommSetup(ncclComm_t comm) {
|
||||
|
||||
NCCLCHECKGOTO(ncclCudaMemcpyAsync(devCommAndChans, &tmpCommAndChans, 1, comm->sharedRes->deviceStream.cudaStream), ret, fail);
|
||||
exit:
|
||||
CUDACHECK(cudaStreamSynchronize(comm->sharedRes->deviceStream.cudaStream));
|
||||
NCCLCHECK(ncclStrongStreamSynchronize(&comm->sharedRes->deviceStream));
|
||||
NCCLCHECK(ncclStrongStreamRelease(ncclCudaGraphNone(), &comm->sharedRes->deviceStream));
|
||||
return ret;
|
||||
fail:
|
||||
|
||||
@@ -163,11 +163,9 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
|
||||
if (conn->connected == 0) {
|
||||
NCCLCHECKGOTO(conn->transportComm->connect(comm, sendData[i] + sendDataOffset++, 1, comm->rank, conn), ret, fail);
|
||||
if (ret == ncclSuccess) {
|
||||
struct ncclDevChannelPeer* addr;
|
||||
conn->connected = 1;
|
||||
/* comm->channels[c].devPeers[sendPeer]->send[connIndex] is a device memory access. */
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr, &comm->channels[c].devPeers[sendPeer], sizeof(struct ncclDevChannelPeer*), cudaMemcpyDeviceToHost, comm->sharedRes->hostStream.cudaStream), ret, fail);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr->send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[sendPeer]->send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
|
||||
} else if (ret == ncclInProgress) {
|
||||
allChannelsConnected = false;
|
||||
}
|
||||
@@ -183,11 +181,9 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
|
||||
if (conn->connected == 0) {
|
||||
NCCLCHECKGOTO(conn->transportComm->connect(comm, recvData[i] + recvDataOffset++, 1, comm->rank, conn), ret, fail);
|
||||
if (ret == ncclSuccess) {
|
||||
struct ncclDevChannelPeer* addr;
|
||||
conn->connected = 1;
|
||||
/* comm->channels[c].devPeers[recvPeer]->recv[connIndex] is a device memory access. */
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr, &comm->channels[c].devPeers[recvPeer], sizeof(struct ncclDevChannelPeer*), cudaMemcpyDeviceToHost, comm->sharedRes->hostStream.cudaStream), ret, fail);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr->recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[recvPeer]->recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
|
||||
} else if (ret == ncclInProgress) {
|
||||
allChannelsConnected = false;
|
||||
}
|
||||
|
||||
@@ -68,7 +68,7 @@ struct userIbDev {
|
||||
uint16_t port_en;
|
||||
};
|
||||
|
||||
#define MAX_IB_DEVS 16
|
||||
#define MAX_IB_DEVS 32
|
||||
struct ncclIbDev ncclIbDevs[MAX_IB_DEVS];
|
||||
struct userIbDev userIbDevs[MAX_IB_DEVS];
|
||||
pthread_mutex_t ncclIbLock = PTHREAD_MUTEX_INITIALIZER;
|
||||
|
||||
@@ -359,12 +359,10 @@ ncclResult_t ncclNvlsSetup(struct ncclComm* comm, struct ncclComm* parent) {
|
||||
peer->send[0].conn.tail = (uint64_t*)(mem + buffSize + memSize / 2);
|
||||
peer->send[0].conn.flags |= NCCL_NVLS_MIN_POLL;
|
||||
|
||||
struct ncclDevChannelPeer* addr;
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr, comm->channels[c].devPeers + nvlsPeer, sizeof(struct ncclDevChannelPeer*), cudaMemcpyDeviceToHost, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr->send[0], &peer->send[0].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr->recv[0], &peer->recv[0].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr->send[1], &peer->send[1].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&addr->recv[1], &peer->recv[1].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[nvlsPeer]->send[0], &peer->send[0].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[nvlsPeer]->recv[0], &peer->recv[0].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[nvlsPeer]->send[1], &peer->send[1].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[nvlsPeer]->recv[1], &peer->recv[1].conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), res, cleanup);
|
||||
|
||||
/*INFO(NCCL_INIT|NCCL_NVLS, "Peer %d Channel %d MC buff %p/%p UC Buff %p/%p",
|
||||
nvlsPeer, c,
|
||||
|
||||
Ссылка в новой задаче
Block a user