Fix GDRDMA read and remove unused files

[ROCm/rccl commit: 9cc3b56166]
2021-02-09 01:34:39 +00:00
@@ -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<int UNROLL, class FUNC, typename T>
-__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<UNROLL, FUNC, T, 1, 1, 1, 1>(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<UNROLL, GATHER_CHUNKSTEPS/GATHER_SLICESTEPS, GATHER_SLICESTEPS, T, 1, 1, 0, FUNC>
-            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<UNROLL, GATHER_CHUNKSTEPS/GATHER_SLICESTEPS, GATHER_SLICESTEPS, T, 1, 1, 0, FUNC>
-              prims(tid, nthreads, &peerNone, &peerSend, NULL, stepSize, channel, comm);
-
-            ssize_t send_offset = chunkOffset;
-            prims.send(thisInput+send_offset, nelem);
-          }
-        }
-      }
-    }
-  }
-}
@@ -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<int UNROLL, class FUNC, typename T>
-__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<UNROLL, FUNC, T, 1, 1, 1, 1>(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<UNROLL, SCATTER_CHUNKSTEPS/SCATTER_SLICESTEPS, SCATTER_SLICESTEPS, T, 1, 1, 0, FUNC>
-            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<UNROLL, SCATTER_CHUNKSTEPS/SCATTER_SLICESTEPS, SCATTER_SLICESTEPS, T, 1, 1, 0, FUNC>
-              prims(tid, nthreads, &peerRecv, &peerNone, NULL, stepSize, channel, comm);
-
-            ssize_t recv_offset = chunkOffset;
-            prims.recv(thisOutput+recv_offset, nelem);
-          }
-        }
-      }
-    }
-  }
-}
@@ -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;