AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Updating RCCL based on NCCL 2.3.7

Browse Source 
- Contains modifications to support AMD hardware
- Adds unit tests


[ROCm/rccl commit: 55a4b22ad7]
This commit is contained in:
Gilbert Lee
2019-05-16 16:16:18 +00:00
parent 3a7974d4c4
commit 60f91f645d
79 changed files with 2678 additions and 399 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/include/common_coll.h
+11 -10
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2016-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -12,16 +13,16 @@
#include "collectives/collectives.h"
static ncclResult_t PointerCheck(const void* pointer, struct ncclComm* comm, const char* ptrname, const char* opname) {
cudaPointerAttributes attr;
cudaError_t err = cudaPointerGetAttributes(&attr, pointer);
if (err != cudaSuccess || attr.devicePointer == NULL) {
hipPointerAttribute_t attr;
hipError_t err = hipPointerGetAttributes(&attr, pointer);
if (err != hipSuccess || attr.devicePointer == NULL) {
WARN("%s : %s is not a valid pointer", opname, ptrname);
return ncclInvalidArgument;
}
#if CUDART_VERSION >= 10000
if (attr.type == cudaMemoryTypeDevice && attr.device != comm->cudaDev) {
if (attr.type == hipMemoryTypeDevice && attr.device != comm->cudaDev) {
#else
if (attr.memoryType == cudaMemoryTypeDevice && attr.device != comm->cudaDev) {
if (attr.memoryType == hipMemoryTypeDevice && attr.device != comm->cudaDev) {
#endif
WARN("%s : %s allocated on device %d mismatchs with NCCL device %d", opname, ptrname, attr.device, comm->cudaDev);
return ncclInvalidArgument;
@@ -99,7 +100,7 @@ static inline void ncclGetCollResource(ncclComm_t comm, size_t nbytes, int* nrin
int llEnforced = 0; /* see if the size falls in the NCCL_LL_THRESHOLD range set by user */
if (comm->llThreshold >= 0) { /* user sets total LL threshold */
if (nbytes > comm->llThreshold) { /* non-LL */
*nthreads = comm->nThreads+1;
*nthreads = comm->nThreads;
*nrings = comm->nRings;
return;
} else {
@@ -129,12 +130,12 @@ static inline void ncclGetCollResource(ncclComm_t comm, size_t nbytes, int* nrin
nr = DIVUP(nbytes, (NCCL_LL_RING_THRESHOLD*ll_max_nthreads*comm->nRanks)); /* else we try the max number of LL threads */
nr = nr == 0 ? 1 : nr > comm->nRings ? comm->nRings : nr;
*ll = nbytes > comm->nRanks*nr*ll_max_nthreads*comm->threadThreshold ? llEnforced : 1;
*nthreads = *ll ? ll_max_nthreads : comm->nThreads+1;
*nthreads = *ll ? ll_max_nthreads : comm->nThreads;
*nrings = *ll ? (int)nr : comm->nRings;
}
static ncclResult_t saveKernel(int coll, const void* sendbuff, void* recvbuff, size_t count,
ncclDataType_t dtype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream, size_t nbytes, int loopFactor) {
ncclDataType_t dtype, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream, size_t nbytes, int loopFactor) {
int llMode, nBlocks, nThreads;
ncclGetCollResource(comm, nbytes, &nBlocks, &nThreads, &llMode);
comm->myParams->blockDim.x = std::max((int)comm->myParams->blockDim.x, nThreads);
@@ -164,7 +165,7 @@ static ncclResult_t saveKernel(int coll, const void* sendbuff, void* recvbuff, s
int opIndex = ring->collFifoTail;
struct ncclColl* c = ring->collectives+opIndex;
volatile uint8_t* activePtr = (volatile uint8_t*)&c->active;
while (activePtr[0] != 0) sched_yield();
while (LOAD(activePtr) != 0) sched_yield();
struct CollectiveArgs* args = &c->args;
args->root = root;
@@ -180,7 +181,7 @@ static ncclResult_t saveKernel(int coll, const void* sendbuff, void* recvbuff, s
c->nThreads = nThreads;
c->funcIndex = FUNC_INDEX(coll, op, dtype, llMode);
c->active = 1;
STORE(&c->active, 1);
opIndex = (opIndex+1)%NCCL_MAX_OPS;
c->nextIndex = opIndex;
ring->collFifoTail = opIndex;
projects/rccl/src/include/core.h
+41 -24
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -16,16 +17,17 @@
#include <algorithm> // std::min/std::max
#include <unistd.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#if CUDART_VERSION < 9000
struct cudaLaunchParams {
void *func;
void (*func)(struct ncclColl);
dim3 gridDim;
dim3 blockDim;
void **args;
struct ncclColl **args;
size_t sharedMem;
cudaStream_t stream;
hipStream_t stream;
};
#endif
@@ -35,10 +37,10 @@ struct cudaLaunchParams {
// Rings / LL tuning
#define NCCL_LL_RING_THRESHOLD 8 // Per thread size before we start increasing nrings
#define NCCL_THREAD_THRESHOLD 64 // Per thread size before we switch to non-LL for Volta and above
#define NCCL_THREAD_THRESHOLD 256 // Per thread size before we switch to non-LL for Volta and above
#define NCCL_THREAD_THRESHOLD_PREVOLTA 32 // Per thread size before we switch to non-LL for pre-Volta archs
#define NCCL_LL_MAX_NTHREADS 256
#define NCCL_LL_MIN_NTHREADS 64
#define NCCL_LL_MIN_NTHREADS 256
#define DIVUP(x, y) \
(((x)+(y)-1)/(y))
@@ -90,9 +92,9 @@ struct ncclConnector {
struct ncclConnInfo conn;
};
#define CACHE_LINE_SIZE 128
#define CACHE_LINE_SIZE 64
#define MEM_ALIGN 4096
#define SIZES_FIFO_SIZE 32
#define SIZES_FIFO_SIZE 16
#define CUDA_IPC_MIN 2097152UL /* 2MiB - not currently used */
#define NCCL_LL_CHUNKS 8
@@ -164,6 +166,8 @@ struct ncclRing {
};
static_assert(sizeof(struct ncclRing) == 0x80*sizeof(int), "ncclRing must have a pow2 size");
#pragma pack(push) /* push current alignment to stack */
#pragma pack(4) /* set alignment to 4 bytes boundary */
/* CollectiveArgs + ncclColl are to be a power of two, currently 64 bytes, */
/* to make sure reads to host from the CUDA kernel are aligned. */
/* Make sure to adjust padding at the end of ncclColl. */
@@ -197,6 +201,7 @@ struct ncclColl {
};
};
static_assert(sizeof(struct ncclColl) == (0x10*sizeof(int)), "ncclColl must have a pow2 size");
#pragma pack(pop) /* restore original alignment from stack */
struct ncclComm {
struct ncclRing rings[MAXRINGS];
@@ -206,9 +211,9 @@ struct ncclComm {
int cudaDev; // my cuda device index
enum { GROUP, PARALLEL } launchMode;
cudaStream_t userStream;
hipStream_t userStream;
bool userStreamSet;
cudaEvent_t doneEvent;
hipEvent_t doneEvent;
bool checkPointers;
// Counter to make sure collectives match (needed for bcast/reduce
@@ -225,7 +230,7 @@ struct ncclComm {
// An internal CUDA stream for NCCL kernel CGMD launches
int groupCudaStream;
cudaStream_t groupStream;
hipStream_t groupStream;
// Device copy of the communicator
struct ncclComm *devComm;
@@ -243,22 +248,31 @@ struct ncclComm {
int* intraCGMode; // Whether we can use CUDA9 CGMD or not
int* intraCC; // Only to check all have the same ComputeCap and disable CGMode if not
struct ncclColl args;
void* argsptr;
struct ncclColl* argsptr;
};
// Convert volatile access to atomic
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_SEQ_CST)
#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST)
#else
#define LOAD(VAR) *(VAR)
#define STORE(DST, SRC) *(DST) = (SRC)
#endif
// Check CUDA calls
#define CUDACHECK(cmd) do { \
cudaError_t e = cmd; \
if( e != cudaSuccess ) { \
WARN("Cuda failure '%s'", cudaGetErrorString(e)); \
hipError_t e = cmd; \
if( e != hipSuccess ) { \
WARN("Cuda failure '%s'", hipGetErrorString(e)); \
return ncclUnhandledCudaError; \
} \
} while(false)
#define CUDACHECKGOTO(cmd, res, label) do { \
cudaError_t e = cmd; \
if( e != cudaSuccess ) { \
WARN("Cuda failure '%s'", cudaGetErrorString(e)); \
hipError_t e = cmd; \
if( e != hipSuccess ) { \
WARN("Cuda failure '%s'", hipGetErrorString(e)); \
res = ncclUnhandledCudaError; \
goto label; \
} \
@@ -327,14 +341,14 @@ int ncclCudaCompCap();
#include <sys/mman.h>
static inline ncclResult_t ncclCudaHostAlloc(void** ptr, void** devPtr, size_t size) {
CUDACHECK(cudaHostAlloc(ptr, size, cudaHostAllocMapped));
CUDACHECK(hipHostMalloc(ptr, size, hipHostMallocMapped));
memset(*ptr, 0, size);
*devPtr = *ptr;
return ncclSuccess;
}
static inline ncclResult_t ncclCudaHostFree(void* ptr) {
CUDACHECK(cudaFreeHost(ptr));
CUDACHECK(hipHostFree(ptr));
return ncclSuccess;
}
@@ -351,15 +365,18 @@ static ncclResult_t ncclCalloc(T** ptr, size_t nelem) {
}
template <typename T>
static ncclResult_t ncclCudaCalloc(T** ptr, size_t nelem) {
CUDACHECK(cudaMalloc(ptr, nelem*sizeof(T)));
CUDACHECK(cudaMemset(*ptr, 0, nelem*sizeof(T)));
static ncclResult_t ncclCudaCalloc(T** ptr, size_t nelem, bool isFineGrain = false) {
if (isFineGrain)
CUDACHECK(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained));
else
CUDACHECK(hipMalloc(ptr, nelem*sizeof(T)));
CUDACHECK(hipMemset(*ptr, 0, nelem*sizeof(T)));
return ncclSuccess;
}
template <typename T>
static ncclResult_t ncclCudaMemcpy(T* dst, T* src, size_t nelem) {
CUDACHECK(cudaMemcpy(dst, src, nelem*sizeof(T), cudaMemcpyDefault));
CUDACHECK(hipMemcpy(dst, src, nelem*sizeof(T), hipMemcpyDefault));
return ncclSuccess;
}
projects/rccl/src/include/debug.h
+1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
projects/rccl/src/include/enqueue.h
+3 -2
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -11,11 +12,11 @@
#include "group.h"
typedef ncclResult_t(*ncclFunc_t)(const void* sendbuff, void* recvbuff, size_t count,
ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream);
ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream);
ncclResult_t ncclEnqueueCheck(ncclFunc_t func, const char* primName, const void* sendbuff,
void* recvbuff, size_t count, ncclDataType_t type, ncclRedOp_t op, int root,
ncclComm_t comm, cudaStream_t stream);
ncclComm_t comm, hipStream_t stream);
ncclResult_t ncclCpuBarrierIn(ncclComm_t comm, int* isLast);
ncclResult_t ncclCpuBarrierLast(ncclComm_t comm);
ncclResult_t ncclCpuBarrierOut(ncclComm_t comm);
projects/rccl/src/include/group.h
+2 -1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2017, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -18,7 +19,7 @@ typedef ncclResult_t(*ncclInitFunc_t)(ncclComm_t* newcomm, int ndev, ncclUniqueI
ncclResult_t ncclAsyncInit(ncclInitFunc_t func, int cudaDev, ncclComm_t* newcomm, int ndev, ncclUniqueId commId, int myrank);
typedef ncclResult_t(*ncclCollFunc_t)(const void* sendbuff, void* recvbuff, size_t count,
ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream);
ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream);
ncclResult_t ncclAsyncColl(ncclComm_t comm);
#endif
projects/rccl/src/include/nvlink.h
+1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2016-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
projects/rccl/src/include/nvlink_stub.h
+20
View File
@@ -0,0 +1,20 @@
/*************************************************************************
* Copyright (c) 2016-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#ifndef NCCL_NVLINK_H_
#define NCCL_NVLINK_H_
#include "topo.h"
#define CONNECT_NVLINK 0x10
#define CONNECT_NVSWITCH 0x100
static int getNumNvlinks(const char* busId) {
return 0;
}
#endif
projects/rccl/src/include/rings.h
+5
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -9,7 +10,11 @@
static int getDefaultThreads() {
// On Kepler, rings are doubled later.
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
return 256;
#else
return ncclCudaCompCap() == 3 ? 128 : 256;
#endif
}
ncclResult_t ncclGetRings(int* nrings, int* nthreads, int rank, int nranks, int* transports, ncclTvalue_t* values, int* prev, int* next);
projects/rccl/src/include/shm.h
+5 -4
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2016-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -39,14 +40,14 @@ static ncclResult_t shmOpen(const char* shmname, const int shmsize, void** shmPt
ncclResult_t res = ncclSuccess;
NCCLCHECKGOTO(shmSetup(shmname, shmsize, &fd, &ptr, create), res, sysError);
CUDACHECKGOTO(cudaHostRegister(ptr, shmsize, cudaHostRegisterMapped), res, cudaError);
CUDACHECKGOTO(cudaHostGetDevicePointer(devShmPtr, ptr, 0), res, cudaError);
CUDACHECKGOTO(hipHostRegister(ptr, shmsize, hipHostRegisterMapped), res, hipError_t);
CUDACHECKGOTO(hipHostGetDevicePointer(devShmPtr, ptr, 0), res, hipError_t);
*shmPtr = ptr;
return ncclSuccess;
sysError:
WARN("Error while %s shared memory segment %s (size %d)\n", create ? "creating" : "attaching to", shmname, shmsize);
cudaError:
hipError_t:
if (fd != -1) close(fd);
if (create) shm_unlink(shmname);
if (ptr != MAP_FAILED) munmap(ptr, shmsize);
@@ -60,7 +61,7 @@ static ncclResult_t shmUnlink(const char* shmname) {
}
static ncclResult_t shmClose(void* shmPtr, void* devShmPtr, const int shmsize) {
CUDACHECK(cudaHostUnregister(shmPtr));
CUDACHECK(hipHostUnregister(shmPtr));
if (munmap(shmPtr, shmsize) != 0) {
WARN("munmap of shared memory failed");
return ncclSystemError;
projects/rccl/src/include/topo.h
+31 -1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2016-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -11,13 +12,41 @@
#include <limits.h>
#include <stdlib.h>
#include <ctype.h>
#include <iostream>
#include <fstream>
#include <string>
#define BUSID_SIZE (sizeof("0000:00:00.0"))
#define BUSID_REDUCED_SIZE (sizeof("0000:00"))
static bool isEPYC() {
std::ifstream cpuinfo("/proc/cpuinfo");
std::string line;
int needed = 2;
static bool vendor_id = true, cpu_family = false, initialized = false;
if (initialized) return (vendor_id && cpu_family);
while (std::getline(cpuinfo, line)) {
if (line.compare(0, 9, "vendor_id") == 0) {
if(line.find("AuthenticAMD") == std::string::npos)
vendor_id = false;
needed --;
}
if (line.compare(0, 10, "cpu family") == 0) {
std::string family_str = line.substr(line.find(": ") + 2);
if (std::stoi(family_str) >= 23)
cpu_family = true;
needed --;
}
if (!needed)
break;
}
initialized = true;
return (vendor_id && cpu_family);
}
static ncclResult_t getCudaPath(int cudaDev, char** path) {
char busId[BUSID_SIZE];
CUDACHECK(cudaDeviceGetPCIBusId(busId, BUSID_SIZE, cudaDev));
CUDACHECK(hipDeviceGetPCIBusId(busId, BUSID_SIZE, cudaDev));
for (int i=0; i<BUSID_SIZE; i++) busId[i] = tolower(busId[i]);
char busPath[] = "/sys/class/pci_bus/0000:00/../../0000:00:00.0";
memcpy(busPath+sizeof("/sys/class/pci_bus/")-1, busId, BUSID_REDUCED_SIZE-1);
@@ -50,6 +79,7 @@ static int pciDistance(char* path1, char* path2) {
if (same == 1) score++;
}
}
if (isEPYC() && score <= 3) return PATH_PHB;
if (score <= 3) return PATH_SOC;
if (score == 4) return PATH_PHB;
if (score == depth-1) return PATH_PIX;