From 020800d7b77ee5376916dffac9881bbcff4fbc36 Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Wed, 2 Aug 2023 16:03:45 +0200 Subject: [PATCH 1/4] Fix inter-node NVLS graph search We were passing a net ID instead of a gpu index, which could cause crashes if those were unrelated (and they usually are). Issue #931 [ROCm/rccl commit: 8ed014bae9f30d4470cdfa655f32d35ee5b3e7ca] --- projects/rccl/src/graph/search.cc | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/graph/search.cc b/projects/rccl/src/graph/search.cc index 3b13c48661..bb83d5d7e3 100644 --- a/projects/rccl/src/graph/search.cc +++ b/projects/rccl/src/graph/search.cc @@ -530,7 +530,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, nets[graph->nChannels], &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 From 26b91b9dbb714b1e58cac04152eec4c035c24db8 Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Wed, 23 Aug 2023 06:30:05 -0700 Subject: [PATCH 2/4] 2.18.5-1 Fix NVLS search (issue #931). Increase max IB NICs to 32. Fix inconsistent device ordering (issue #820). Try to use different devices for different GPUs in systems with more than one NIC per GFU. [ROCm/rccl commit: 559b70f86c190a0d8f67f0d7a0f2c9810dd1e8c7] --- projects/rccl/makefiles/version.mk | 2 +- projects/rccl/src/graph/paths.cc | 4 +- projects/rccl/src/graph/search.cc | 32 +++++-- projects/rccl/src/graph/topo.cc | 133 +++++++------------------- projects/rccl/src/graph/topo.h | 10 ++ projects/rccl/src/graph/xml.cc | 13 ++- projects/rccl/src/transport/net_ib.cc | 2 +- 7 files changed, 80 insertions(+), 116 deletions(-) diff --git a/projects/rccl/makefiles/version.mk b/projects/rccl/makefiles/version.mk index ba162237d4..df96037276 100644 --- a/projects/rccl/makefiles/version.mk +++ b/projects/rccl/makefiles/version.mk @@ -1,6 +1,6 @@ ##### version NCCL_MAJOR := 2 NCCL_MINOR := 18 -NCCL_PATCH := 3 +NCCL_PATCH := 5 NCCL_SUFFIX := PKG_REVISION := 1 diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc index 29cea27f22..450ba658f1 100644 --- a/projects/rccl/src/graph/paths.cc +++ b/projects/rccl/src/graph/paths.cc @@ -730,9 +730,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; cp2pnChannels; c++) { - int mirror = 0; - for (int b=1, mb=(comm->p2pnChannels>>1); bp2pnChannels; b<<=1, mb>>=1) if (c & b) mirror |= mb; - comm->p2pChannels[c] = mirror; + comm->p2pChannels[c] = mirrorBits(c, comm->p2pnChannels); } return ncclSuccess; } diff --git a/projects/rccl/src/graph/search.cc b/projects/rccl/src/graph/search.cc index bb83d5d7e3..dd8896bd20 100644 --- a/projects/rccl/src/graph/search.cc +++ b/projects/rccl/src/graph/search.cc @@ -376,6 +376,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; gnodes[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; inodes[GPU].count; g++) { if (gpu != -1 && gpu != g) continue; @@ -385,14 +407,6 @@ ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, in for (int n=0; nnodes[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; rnodes[GPU].nodes[g].gpu.dev % localNetCount; r++) { - int net0 = localNets[0]; - for (int i=0; ipattern == NCCL_TOPO_PATTERN_NVLS) { if (graph->nChannels < netcount) { int gpu; - NCCLCHECK(ncclTopoGetLocalGpu(system, nets[graph->nChannels], &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 { diff --git a/projects/rccl/src/graph/topo.cc b/projects/rccl/src/graph/topo.cc index 529eff4a41..cdcc0664f7 100644 --- a/projects/rccl/src/graph/topo.cc +++ b/projects/rccl/src/graph/topo.cc @@ -679,126 +679,57 @@ 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; nnodes[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; inodes[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= 64) return ncclInternalError; - *localNetMask |= 1ULL<nodes[GPU].count; - NCCLCHECK(ncclCalloc(&localNetMasks, ngpus)); - - // Fill localNetMasks for all GPUs. - for (int g=0; gnodes[GPU].nodes[gpu].gpu.dev; + if (isPow2(localNetCount)) net = mirrorBits(net, localNetCount); + 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; gnodes[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; } /****************************/ diff --git a/projects/rccl/src/graph/topo.h b/projects/rccl/src/graph/topo.h index 281d16ad9c..8951505fd6 100644 --- a/projects/rccl/src/graph/topo.h +++ b/projects/rccl/src/graph/topo.h @@ -208,4 +208,14 @@ static float ncclTopoNVLinkBw(int cudaCompCap) { cudaCompCap >= 60 ? SM60_NVLINK_BW : SM80_NVLINK_BW; } + +// 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>=1) if (val & b) mirror |= mb; + return mirror; +} #endif diff --git a/projects/rccl/src/graph/xml.cc b/projects/rccl/src/graph/xml.cc index d441de02df..ac862a4e06 100644 --- a/projects/rccl/src/graph/xml.cc +++ b/projects/rccl/src/graph/xml.cc @@ -576,7 +576,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; snSubs; 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)); diff --git a/projects/rccl/src/transport/net_ib.cc b/projects/rccl/src/transport/net_ib.cc index 861fa57f0b..421f0a13a1 100644 --- a/projects/rccl/src/transport/net_ib.cc +++ b/projects/rccl/src/transport/net_ib.cc @@ -66,7 +66,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; From d03ae00bac882f80cff5dd397f4ed0d1a4f17fbf Mon Sep 17 00:00:00 2001 From: Kaiming Ouyang Date: Wed, 20 Sep 2023 05:51:14 -0700 Subject: [PATCH 3/4] Fix cudaMemcpyAsync bug We are trying to use the copy result of first cudaMemcpyAsync in the second cudaMemcpyAsync without sync in between. This patch fixes it by allocating a CPU side array to cache device side addr so that we can avoid this consecutive cuda mem copy. Fixes #957 [ROCm/rccl commit: 4365458757e4107ecbf629b2fd6e0e19a5d237c2] --- projects/rccl/src/channel.cc | 12 ++++++++++++ projects/rccl/src/include/comm.h | 2 ++ projects/rccl/src/init.cc | 2 +- projects/rccl/src/transport.cc | 8 ++------ projects/rccl/src/transport/nvls.cc | 10 ++++------ 5 files changed, 21 insertions(+), 13 deletions(-) diff --git a/projects/rccl/src/channel.cc b/projects/rccl/src/channel.cc index 3edcc2fc2d..245dfd58cb 100644 --- a/projects/rccl/src/channel.cc +++ b/projects/rccl/src/channel.cc @@ -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,6 +54,8 @@ 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)); return ncclSuccess; @@ -77,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 { @@ -86,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; @@ -114,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)); @@ -121,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; @@ -156,5 +166,7 @@ ncclResult_t freeChannel(struct ncclChannel* channel, int nRanks, int collnetNRa } } } + + free(channel->devPeersHostPtr); return ncclSuccess; } diff --git a/projects/rccl/src/include/comm.h b/projects/rccl/src/include/comm.h index e79bf549ba..8986f9349c 100644 --- a/projects/rccl/src/include/comm.h +++ b/projects/rccl/src/include/comm.h @@ -124,6 +124,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; diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 1ea1d7efa3..309ce10bb9 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -437,7 +437,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: diff --git a/projects/rccl/src/transport.cc b/projects/rccl/src/transport.cc index f4b8a2a677..9817beb183 100644 --- a/projects/rccl/src/transport.cc +++ b/projects/rccl/src/transport.cc @@ -147,11 +147,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; } @@ -167,11 +165,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; } diff --git a/projects/rccl/src/transport/nvls.cc b/projects/rccl/src/transport/nvls.cc index 633cb04d8e..07be99d9fe 100644 --- a/projects/rccl/src/transport/nvls.cc +++ b/projects/rccl/src/transport/nvls.cc @@ -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, From 90d7d1b07b66e1f34e902e06fc6e7bb7ebfb9425 Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Tue, 10 Oct 2023 00:34:18 -0700 Subject: [PATCH 4/4] 2.18.6-1 [ROCm/rccl commit: 0b083e52096c387bad7a5c5c65b26a9dca54de8c] --- projects/rccl/makefiles/version.mk | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rccl/makefiles/version.mk b/projects/rccl/makefiles/version.mk index df96037276..fde92c08a0 100644 --- a/projects/rccl/makefiles/version.mk +++ b/projects/rccl/makefiles/version.mk @@ -1,6 +1,6 @@ ##### version NCCL_MAJOR := 2 NCCL_MINOR := 18 -NCCL_PATCH := 5 +NCCL_PATCH := 6 NCCL_SUFFIX := PKG_REVISION := 1