diff options
Diffstat (limited to 'src/channel.cc')
-rw-r--r-- | src/channel.cc | 56 |
1 files changed, 56 insertions, 0 deletions
diff --git a/src/channel.cc b/src/channel.cc new file mode 100644 index 0000000..b053e5b --- /dev/null +++ b/src/channel.cc @@ -0,0 +1,56 @@ +/************************************************************************* + * Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +#include "channel.h" +#include "param.h" + +NCCL_PARAM(Buffsize, "BUFFSIZE", DEFAULT_BUFFER_SIZE_BYTES); + +ncclResult_t initChannel(struct ncclComm* comm, int channelid) { + struct ncclChannel* channel = comm->channels+channelid; + channel->id = channelid; + + // Setup intermediate buffering + channel->buffSize = ncclParamBuffsize(); + + // Ring index to user rank table. + NCCLCHECK(ncclCudaCalloc(&channel->ring.devUserRanks, comm->nRanks)); + NCCLCHECK(ncclCalloc(&channel->ring.userRanks, comm->nRanks)); + + // Communication structures with peers. + NCCLCHECK(ncclCudaCalloc(&channel->devPeers, comm->nRanks)); + NCCLCHECK(ncclCalloc(&channel->peers, comm->nRanks)); + for (size_t i=0; i<comm->nRanks; ++i) { + channel->peers[i].send.comm = comm; + channel->peers[i].recv.comm = comm; + } + + // Per-channel operation list. + NCCLCHECK(ncclCudaHostAlloc((void**)&channel->collectives, (void**)&channel->devCollectives, sizeof(struct ncclColl)*NCCL_MAX_OPS)); + return ncclSuccess; +} + +ncclResult_t freeChannel(struct ncclChannel* channel, int nRanks) { + // Operation list + NCCLCHECK(ncclCudaHostFree(channel->collectives)); + + // Free Ring index to rank tables + free(channel->ring.userRanks); + CUDACHECK(cudaFree(channel->ring.devUserRanks)); + + // Free transport proxy resources + for (int r=0; r<nRanks; r++) { + struct ncclPeer* peer = channel->peers+r; + if (peer->send.transportResources) NCCLCHECK(peer->send.transportComm->free(peer->send.transportResources)); + if (peer->recv.transportResources) NCCLCHECK(peer->recv.transportComm->free(peer->recv.transportResources)); + } + + // Free the peer structures. + CUDACHECK(cudaFree(channel->devPeers)); + free(channel->peers); + + return ncclSuccess; +} |