From 602fe60428f1e2c0ad4cb31c4dacc80361d0fb5c Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Tue, 9 Feb 2021 01:34:39 +0000 Subject: [PATCH] Fix GDRDMA read and remove unused files [ROCm/rccl commit: 9cc3b5616601c9cc20ae7d93461e4416fa713cc4] --- projects/rccl/src/collectives/device/gather.h | 75 ------------------- .../rccl/src/collectives/device/scatter.h | 75 ------------------- projects/rccl/src/transport/net.cc | 46 ++++++------ 3 files changed, 23 insertions(+), 173 deletions(-) delete mode 100644 projects/rccl/src/collectives/device/gather.h delete mode 100644 projects/rccl/src/collectives/device/scatter.h diff --git a/projects/rccl/src/collectives/device/gather.h b/projects/rccl/src/collectives/device/gather.h deleted file mode 100644 index b6c3865875..0000000000 --- a/projects/rccl/src/collectives/device/gather.h +++ /dev/null @@ -1,75 +0,0 @@ -/************************************************************************* - * Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved. - * Modifications Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ - -#include "devcomm.h" -#include "primitives.h" -#include "collectives.h" - -template -__attribute__((noinline)) -__device__ void ncclGatherKernel(struct CollectiveArgs* args) { - const int tid = threadIdx.x; - const int nthreads = args->coll.nThreads; - const int nChannels = args->coll.nChannels; - struct ncclDevComm* comm = args->comm; - struct ncclChannel* channel = comm->channels+blockIdx.x; - struct ncclRing* ring = &channel->ring; - const ssize_t size = args->coll.count; - const int nranks = comm->nRanks; - const int bid = args->coll.bid; - const int rank = ring->devUserRanks[0]; - const int stepSize = comm->buffSizes[NCCL_PROTO_SIMPLE] / (sizeof(T)*NCCL_STEPS); - const int chunkSize = stepSize * GATHER_CHUNKSTEPS; - const int peersPerChan = DIVUP(nranks, nChannels); - const ssize_t loopSize = (peersPerChan == 1 ? (nChannels/nranks)*(ssize_t)chunkSize : (ssize_t)chunkSize); - const int root = args->coll.root; - - // Compute pointers - const T * __restrict__ thisInput = (const T*)args->sendbuff; - T * __restrict__ thisOutput = (T*)args->recvbuff; - - for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) { - for (int i = 0; i < peersPerChan; i++) { - if ((peersPerChan == 1 && blockIdx.x >= (nChannels/nranks)*nranks) || - (peersPerChan > 1 && blockIdx.x*peersPerChan+i >= nranks)) - continue; - int realChunkSize = min(chunkSize, DIVUP(size-gridOffset, (peersPerChan == 1 ? (nChannels/nranks) : 1))); - ALIGN_SIZE(realChunkSize, nthreads*sizeof(uint64_t)/sizeof(T)); - ssize_t chunkOffset = gridOffset + (peersPerChan == 1 ? (bid/nranks)*realChunkSize : 0); - int nelem = min(realChunkSize, size-chunkOffset); - if ((blockIdx.x*peersPerChan+i)%nranks == 0 && rank == root) { - const T* sendbuff = thisInput+chunkOffset; - T* recvbuff = thisOutput+chunkOffset+rank*size; - if (tid < nthreads && sendbuff != recvbuff) { - // local copy - ReduceOrCopyMulti(tid, nthreads, 1, &sendbuff, 1, &recvbuff, nelem); - } - } - else { - int peerSend = (rank+(blockIdx.x*peersPerChan)+i)%nranks; - int peerRecv = (2*nranks+rank-((blockIdx.x*peersPerChan)%nranks)-(i%nranks))%nranks; - int peerNone = -1; - if (rank == root) { - ncclPrimitives - prims(tid, nthreads, &peerRecv, &peerNone, NULL, stepSize, channel, comm); - - ssize_t recv_offset = chunkOffset + peerRecv*size; - prims.recv(thisOutput+recv_offset, nelem); - } - else { - if (peerSend == root) { - ncclPrimitives - prims(tid, nthreads, &peerNone, &peerSend, NULL, stepSize, channel, comm); - - ssize_t send_offset = chunkOffset; - prims.send(thisInput+send_offset, nelem); - } - } - } - } - } -} diff --git a/projects/rccl/src/collectives/device/scatter.h b/projects/rccl/src/collectives/device/scatter.h deleted file mode 100644 index 848b47713d..0000000000 --- a/projects/rccl/src/collectives/device/scatter.h +++ /dev/null @@ -1,75 +0,0 @@ -/************************************************************************* - * Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved. - * Modifications Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ - -#include "devcomm.h" -#include "primitives.h" -#include "collectives.h" - -template -__attribute__((noinline)) -__device__ void ncclScatterKernel(struct CollectiveArgs* args) { - const int tid = threadIdx.x; - const int nthreads = args->coll.nThreads; - const int nChannels = args->coll.nChannels; - struct ncclDevComm* comm = args->comm; - struct ncclChannel* channel = comm->channels+blockIdx.x; - struct ncclRing* ring = &channel->ring; - const ssize_t size = args->coll.count; - const int nranks = comm->nRanks; - const int bid = args->coll.bid; - const int rank = ring->devUserRanks[0]; - const int stepSize = comm->buffSizes[NCCL_PROTO_SIMPLE] / (sizeof(T)*NCCL_STEPS); - const int chunkSize = stepSize * SCATTER_CHUNKSTEPS; - const int peersPerChan = DIVUP(nranks, nChannels); - const ssize_t loopSize = (peersPerChan == 1 ? (nChannels/nranks)*(ssize_t)chunkSize : (ssize_t)chunkSize); - const int root = args->coll.root; - - // Compute pointers - const T * __restrict__ thisInput = (const T*)args->sendbuff; - T * __restrict__ thisOutput = (T*)args->recvbuff; - - for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) { - for (int i = 0; i < peersPerChan; i++) { - if ((peersPerChan == 1 && blockIdx.x >= (nChannels/nranks)*nranks) || - (peersPerChan > 1 && blockIdx.x*peersPerChan+i >= nranks)) - continue; - int realChunkSize = min(chunkSize, DIVUP(size-gridOffset, (peersPerChan == 1 ? (nChannels/nranks) : 1))); - ALIGN_SIZE(realChunkSize, nthreads*sizeof(uint64_t)/sizeof(T)); - ssize_t chunkOffset = gridOffset + (peersPerChan == 1 ? (bid/nranks)*realChunkSize : 0); - int nelem = min(realChunkSize, size-chunkOffset); - if ((blockIdx.x*peersPerChan+i)%nranks == 0 && rank == root) { - const T* sendbuff = thisInput+chunkOffset+rank*size; - T* recvbuff = thisOutput+chunkOffset; - if (tid < nthreads && sendbuff != recvbuff) { - // local copy - ReduceOrCopyMulti(tid, nthreads, 1, &sendbuff, 1, &recvbuff, nelem); - } - } - else { - int peerSend = (rank+(blockIdx.x*peersPerChan)+i)%nranks; - int peerRecv = (2*nranks+rank-((blockIdx.x*peersPerChan)%nranks)-(i%nranks))%nranks; - int peerNone = -1; - if (rank == root) { - ncclPrimitives - prims(tid, nthreads, &peerNone, &peerSend, NULL, stepSize, channel, comm); - - ssize_t send_offset = chunkOffset + peerSend*size; - prims.send(thisInput+send_offset, nelem); - } - else { - if (peerRecv == root) { - ncclPrimitives - prims(tid, nthreads, &peerRecv, &peerNone, NULL, stepSize, channel, comm); - - ssize_t recv_offset = chunkOffset; - prims.recv(thisOutput+recv_offset, nelem); - } - } - } - } - } -} diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc index 1f04d3ef13..84f3d3d4d1 100644 --- a/projects/rccl/src/transport/net.cc +++ b/projects/rccl/src/transport/net.cc @@ -101,29 +101,6 @@ ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st if (resources->buffSizes[LOC_HOSTMEM]) { NCCLCHECK(ncclCudaHostCalloc(resources->buffers+LOC_HOSTMEM, resources->buffSizes[LOC_HOSTMEM])); } - if (resources->useGdr) { - //CUDACHECK(hipDeviceGetAttribute((int*)&resources->curr_hdp_reg, hipDeviceAttributeHdpMemFlushCntl, myInfo->cudaDev)); - struct data_struct {hsa_agent_t agent; int counter;} out; - out.counter = 0; - out.agent.handle = myInfo->cudaDev; - hsa_iterate_agents([](hsa_agent_t agent, void* data) { - int devId = ((struct data_struct *)data)->agent.handle; - hsa_device_type_t type; - hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); - if(type != HSA_DEVICE_TYPE_GPU) - return HSA_STATUS_SUCCESS; - if(((struct data_struct *)data)->counter!=devId) { - ((struct data_struct *)data)->counter++; - return HSA_STATUS_SUCCESS; - } - ((struct data_struct *)data)->agent = agent; - return HSA_STATUS_SUCCESS; - }, (void*)&out); - hsa_amd_hdp_flush_t hdpinfo; - hsa_status_t err = hsa_agent_get_info(out.agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_HDP_FLUSH, &hdpinfo); - resources->curr_hdp_reg = hdpinfo.HDP_MEM_FLUSH_CNTL; - send->conn.curr_hdp_reg = resources->curr_hdp_reg; - } int offsets[LOC_COUNT]; offsets[LOC_HOSTMEM] = offsets[LOC_DEVMEM] = 0; @@ -134,6 +111,29 @@ ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st } } + if (resources->useGdr) { + //CUDACHECK(hipDeviceGetAttribute((int*)&resources->curr_hdp_reg, hipDeviceAttributeHdpMemFlushCntl, myInfo->cudaDev)); + struct data_struct {hsa_agent_t agent; int counter;} out; + out.counter = 0; + out.agent.handle = myInfo->cudaDev; + hsa_iterate_agents([](hsa_agent_t agent, void* data) { + int devId = ((struct data_struct *)data)->agent.handle; + hsa_device_type_t type; + hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); + if(type != HSA_DEVICE_TYPE_GPU) + return HSA_STATUS_SUCCESS; + if(((struct data_struct *)data)->counter!=devId) { + ((struct data_struct *)data)->counter++; + return HSA_STATUS_SUCCESS; + } + ((struct data_struct *)data)->agent = agent; + return HSA_STATUS_SUCCESS; + }, (void*)&out); + hsa_amd_hdp_flush_t hdpinfo; + hsa_status_t err = hsa_agent_get_info(out.agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_HDP_FLUSH, &hdpinfo); + resources->curr_hdp_reg = hdpinfo.HDP_MEM_FLUSH_CNTL; + send->conn.curr_hdp_reg = resources->curr_hdp_reg; + } INFO(NCCL_INIT|NCCL_NET,"Channel %02d : %d[%lx] -> %d[%lx] [send] via NET/%s/%d%s%s", channelId, myInfo->rank, myInfo->busId, peerInfo->rank, peerInfo->busId, ncclNetName(), resources->netDev, resources->useGdr ? "/GDRDMA" : "", resources->shared ? "/Shared" : ""); return ncclSuccess;