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

RCCL 2.4 update

Browse Source
This commit is contained in:
Wenkai Du
2019-07-05 15:43:00 -07:00
parent 4d579e51cc
commit f11c8f60cd
95 changed files with 7829 additions and 614 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/collectives/device/all_gather.cu
+6 -3
View File
@@ -1,11 +1,14 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* 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
************************************************************************/
#include "all_gather.h"
#include "common.h"
#include "all_gather.h"
#include "collectives.h"
IMPL_COLL_C(ncclAllGather, ncclCollAllGather);
#define UNROLL 4
IMPL_COLL3(ncclAllGather, copy, FuncSum, i8, int8_t, ncclCollAllGather, ncclSum, ncclInt8);
src/collectives/device/all_gather.h
+6 -1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -9,9 +10,10 @@
#include "collectives.h"
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllGatherRingKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int nthreads = blockDim.x - 1;
const int nthreads = blockDim.x;
const int bid = args->bid;
struct ncclDevComm* comm = args->comm;
struct ncclChannel* channel = comm->channels+blockIdx.x;
@@ -67,9 +69,11 @@ __device__ void ncclAllGatherRingKernel(struct CollectiveArgs* args) {
}
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllGatherTreeKernel(struct CollectiveArgs* args) { }
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllGatherRingLLKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int bid = args->bid;
@@ -128,4 +132,5 @@ __device__ void ncclAllGatherRingLLKernel(struct CollectiveArgs* args) {
}
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllGatherTreeLLKernel(struct CollectiveArgs* args) { }
src/collectives/device/all_reduce.cu
+9 -3
View File
@@ -1,11 +1,17 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* 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
************************************************************************/
#include "all_reduce.h"
#include "common.h"
#include "all_reduce.h"
#include "collectives.h"
IMPL_COLL_R(ncclAllReduce, ncclCollAllReduce);
#define UNROLL 4
IMPL_COLL2(ncclAllReduce, sum, FuncSum, ncclCollAllReduce, ncclSum);
IMPL_COLL2(ncclAllReduce, prod, FuncProd, ncclCollAllReduce, ncclProd);
IMPL_COLL2(ncclAllReduce, min, FuncMin, ncclCollAllReduce, ncclMin);
IMPL_COLL2(ncclAllReduce, max, FuncMax, ncclCollAllReduce, ncclMax);
src/collectives/device/all_reduce.h
+25 -2
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -9,9 +10,10 @@
#include "collectives.h"
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int nthreads = blockDim.x - 1;
const int nthreads = blockDim.x;
const int bid = args->bid;
struct ncclDevComm* comm = args->comm;
struct ncclChannel* channel = comm->channels+blockIdx.x;
@@ -21,6 +23,11 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
const int stepSize = channel->buffSize / (sizeof(T)*NCCL_STEPS);
const int chunkSize = stepSize * ALLREDUCE_CHUNKSTEPS;
const ssize_t loopSize = args->nChannels*(ssize_t)chunkSize;
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t clk, t0 = 0ULL, ws, wr;
if (tid == 0) clk = clock64();
#endif
// Compute pointers
const T * __restrict__ thisInput = (const T*)args->ThisInput;
@@ -44,7 +51,9 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
offset = chunkOffset + slice * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
prims.send(thisInput+offset, nelem);
ACCUMULATE_COUNTER(send);
// k-2 steps: reduce and copy to next GPU
for (int j=2; j<nranks; ++j) {
@@ -52,7 +61,9 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
offset = chunkOffset + slice * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
prims.recvReduceSend(thisInput+offset, nelem);
ACCUMULATE_COUNTER(recvReduceSend);
}
// step k-1: reduce this buffer and data, which will produce the final
@@ -61,7 +72,9 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
offset = chunkOffset + slice * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
prims.directRecvReduceCopySend(thisInput+offset, thisOutput+offset, offset, nelem);
ACCUMULATE_COUNTER(directRecvReduceCopySend);
// k-2 steps: copy to next GPU
for (int j=1; j<nranks-1; ++j) {
@@ -69,7 +82,9 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
offset = chunkOffset + slice * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
prims.directRecvCopySend(thisOutput+offset, offset, nelem);
ACCUMULATE_COUNTER(directRecvCopySend);
}
// Make final copy from buffer to dest.
@@ -78,14 +93,20 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
nelem = min(realChunkSize, size-offset);
// Final wait/copy.
INIT_COUNTER;
prims.directRecv(thisOutput+offset, offset, nelem);
ACCUMULATE_COUNTER(directRecv);
}
#ifdef ENABLE_PROFILING
if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), clock64() - clk, __ATOMIC_SEQ_CST);
#endif
}
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllReduceTreeKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int nthreads = blockDim.x - 1;
const int nthreads = blockDim.x;
const int bid = args->bid;
struct ncclDevComm* comm = args->comm;
struct ncclChannel* channel = comm->channels+blockIdx.x;
@@ -135,6 +156,7 @@ __device__ void ncclAllReduceTreeKernel(struct CollectiveArgs* args) {
}
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllReduceRingLLKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int bid = args->bid;
@@ -210,6 +232,7 @@ __device__ void ncclAllReduceRingLLKernel(struct CollectiveArgs* args) {
}
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclAllReduceTreeLLKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int nthreads = args->nThreads;
src/collectives/device/broadcast.cu
+6 -3
View File
@@ -1,11 +1,14 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* 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
************************************************************************/
#include "broadcast.h"
#include "common.h"
#include "broadcast.h"
#include "collectives.h"
IMPL_COLL_C(ncclBroadcast, ncclCollBroadcast);
#define UNROLL 4
IMPL_COLL3(ncclBroadcast, copy, FuncSum, i8, int8_t, ncclCollBroadcast, ncclSum, ncclInt8);
src/collectives/device/broadcast.h
+22 -1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -9,9 +10,10 @@
#include "collectives.h"
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclBroadcastRingKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int nthreads = blockDim.x - 1;
const int nthreads = blockDim.x;
const int bid = args->bid;
struct ncclDevComm* comm = args->comm;
struct ncclChannel* channel = comm->channels+blockIdx.x;
@@ -23,6 +25,11 @@ __device__ void ncclBroadcastRingKernel(struct CollectiveArgs* args) {
const int rank = ring->devUserRanks[0];
const int nextRank = ring->devUserRanks[1];
const int root = args->root;
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t clk, t0 = 0ULL, ws, wr;
if (tid == 0) clk = clock64();
#endif
// Compute pointers
const T * __restrict__ thisInput = (const T*)args->ThisInput;
@@ -39,22 +46,35 @@ __device__ void ncclBroadcastRingKernel(struct CollectiveArgs* args) {
if (rank == root) {
if (thisInput == thisOutput) {
INIT_COUNTER;
prims.send(thisInput+offset, nelem);
ACCUMULATE_COUNTER(send);
} else {
INIT_COUNTER;
prims.copySend(thisInput+offset, thisOutput+offset, nelem);
ACCUMULATE_COUNTER(copySend);
}
} else if (nextRank == root) {
INIT_COUNTER;
prims.recv(thisOutput+offset, nelem);
ACCUMULATE_COUNTER(recv);
} else {
INIT_COUNTER;
prims.recvCopySend(thisOutput+offset, nelem);
ACCUMULATE_COUNTER(recvCopySend);
}
}
#ifdef ENABLE_PROFILING
if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), clock64() - clk, __ATOMIC_SEQ_CST);
#endif
}
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclBroadcastTreeKernel(struct CollectiveArgs* args) { }
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclBroadcastRingLLKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int bid = args->bid;
@@ -99,4 +119,5 @@ __device__ void ncclBroadcastRingLLKernel(struct CollectiveArgs* args) {
}
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclBroadcastTreeLLKernel(struct CollectiveArgs* args) { }
src/collectives/device/common.h
+127 -61
View File
@@ -1,5 +1,7 @@
#include "hip/hip_runtime.h"
/*************************************************************************
* Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -10,11 +12,18 @@
#include "../collectives.h"
#include "devcomm.h"
#include "nccl.h"
#include <type_traits>
// Exit If Abort Barrier across CTA: make sure all threads exit consistently
// Each thread sets a predicate to true if abort == 1
// all CTA's threads enter the barrier and do a popc on their predicates being True
// If any of the thread's predicate was True, all the threads call exit()
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#define exitIfAbortBarrier(abort, abortCount) \
if (abort) __atomic_fetch_add(abortCount, 1, __ATOMIC_SEQ_CST); \
__syncthreads(); \
if (LOAD(abortCount)) { asm volatile ("s_endpgm"); return; }
#else
static inline __device__ void exitIfAbortBarrier(int abort) {
uint32_t popc;
asm ("{");
@@ -24,21 +33,116 @@ static inline __device__ void exitIfAbortBarrier(int abort) {
asm ("}");
if (popc) { asm volatile ("exit;"); }
}
#endif
typedef void(*ncclKern_t)(struct CollectiveArgs* args);
extern __device__ ncclKern_t ncclFuncs[];
#define NCCL_FUNC5(coll, op, dtype) \
NCCL_COLL_NAME(coll, op, dtype), \
NCCL_COLL_NAME(coll##LL, op, dtype)
static __device__ void load_parallel(void* dst, void* src, size_t size, int tid) {
#define NCCL_FUNC4(coll, op, dtype) \
NCCL_FUNC5(coll##Ring, op, dtype)
// Must be consistent with ncclDataType_t
#define NCCL_FUNCS3A(coll, op) \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, u8), \
NCCL_FUNC4(coll, op, i32), \
NCCL_FUNC4(coll, op, u32), \
NCCL_FUNC4(coll, op, i64), \
NCCL_FUNC4(coll, op, u64), \
NCCL_FUNC4(coll, op, f16), \
NCCL_FUNC4(coll, op, f32), \
NCCL_FUNC4(coll, op, f64)
#define NCCL_FUNCS3B(coll, op) \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8)
// Must be consistent with ncclRedOp_t
#define NCCL_FUNCS2A(coll) \
NCCL_FUNCS3A(coll, sum ), \
NCCL_FUNCS3A(coll, prod), \
NCCL_FUNCS3A(coll, max ), \
NCCL_FUNCS3A(coll, min )
#define NCCL_FUNCS2B(coll) \
NCCL_FUNCS3B(coll, copy), \
NCCL_FUNCS3B(coll, copy), \
NCCL_FUNCS3B(coll, copy), \
NCCL_FUNCS3B(coll, copy)
// Must be consistent with ncclColl_t
#define NCCL_FUNCS() { \
NCCL_FUNCS2B(ncclBroadcast), \
NCCL_FUNCS2A(ncclReduce), \
NCCL_FUNCS2B(ncclAllGather), \
NCCL_FUNCS2A(ncclReduceScatter), \
NCCL_FUNCS2A(ncclAllReduce) }
// Must be consistent with the ncclFuncSet enum
using ncclFunc_t = void (*)(struct CollectiveArgs*);
static const __device__ constexpr ncclFunc_t ncclFuncs[]{
// Don't try to initialize the host shadow copy of this device-side global
// variable. There is no host pointer to a device-side function, which
// confuses clang. This will be fixed in the next clang release.
#if defined(__HIP_DEVICE_COMPILE__)
NCCL_FUNCS2B(ncclBroadcast),
NCCL_FUNCS2A(ncclReduce),
NCCL_FUNCS2B(ncclAllGather),
NCCL_FUNCS2A(ncclReduceScatter),
NCCL_FUNCS2A(ncclAllReduce)
#endif
};
template<unsigned short f, unsigned short l>
struct Caller {
static
void call(ncclColl* const c) noexcept
{
constexpr unsigned short m = f + (l - f) / 2;
return (c->funcIndex < m) ? Caller<f, m>::call(c) : Caller<m, l>::call(c);
}
};
template<unsigned short f>
struct Caller<f, f + 1>{
static
void call(struct ncclColl* const c) noexcept { ncclFuncs[f](&c->args); }
};
inline
__device__
void NCCL_CALL_FUNCTIONS(struct ncclColl* const c) noexcept {
if (c->funcIndex < 72) {
if (c->funcIndex % 2) ncclBroadcastRingLL_copy_i8(&c->args);
else ncclBroadcastRing_copy_i8(&c->args);
}
else if (c->funcIndex < 144) Caller<72, 144>::call(c);
else if (c->funcIndex < 216) {
if (c->funcIndex % 2) ncclAllGatherRingLL_copy_i8(&c->args);
else ncclAllGatherRing_copy_i8(&c->args);
}
else Caller<216, 360>::call(c);
}
static __device__ void load_parallel(void* dst, void* src, size_t size, int tid, uint32_t* abortCount) {
int* d = (int*)dst;
int* s = (int*)src;
// When aggregation is effective, if some threads have aborted inside the LL kernel,
// make sure the rest of the threads abort as well
exitIfAbortBarrier(0);
exitIfAbortBarrier(0, abortCount);
for (int o = tid; o < (size/sizeof(int)); o += blockDim.x) d[o] = s[o];
__syncthreads();
}
static __device__ void load_coll(struct ncclColl* localColl, struct ncclColl* hostColl, int tid) {
load_parallel(localColl, hostColl, sizeof(struct ncclColl), tid);
static __device__ void load_coll(struct ncclColl* localColl, struct ncclColl* hostColl, int tid, uint32_t* abortCount) {
load_parallel(localColl, hostColl, sizeof(struct ncclColl), tid, abortCount);
if (tid == 0) hostColl->active = 0;
}
@@ -56,23 +160,27 @@ __global__ void NCCL_KERN_NAME(coll, op, dtype)(struct ncclColl firstColl) { \
int tid = threadIdx.x; \
int bid = blockIdx.x; \
__shared__ struct ncclColl localColl; \
__shared__ uint32_t abortCount; \
if (tid == 0) abortCount = 0; \
__syncthreads(); \
\
struct ncclDevComm* comm = firstColl.args.comm; \
struct ncclChannel* channel = comm->channels+bid; \
struct ncclColl* c; \
channel->abortCount = &abortCount; \
if (bid == 0) { \
/* To optimize for latency, (only) the first operation is passed as argument.*/ \
c = &firstColl; \
} else { \
c = &localColl; \
load_coll(c, channel->devCollectives+channel->collFifoHead, tid); \
load_coll(c, channel->devCollectives+channel->collFifoHead, tid, &abortCount); \
} \
while (1) { \
if (tid < c->args.nThreads) { \
if (c->funcIndex == fIndex) { \
coll##Kernel<COLL_UNROLL, ncclFunc<ctype>, ctype>(&c->args); \
} else { \
ncclFuncs[c->funcIndex](&c->args); \
NCCL_CALL_FUNCTIONS(c); \
} \
} \
int nextIndex = c->nextIndex; \
@@ -84,7 +192,7 @@ __global__ void NCCL_KERN_NAME(coll, op, dtype)(struct ncclColl firstColl) { \
\
/* Load next collective operation*/ \
c = &localColl; /* for bid 0 */ \
load_coll(c, channel->devCollectives+nextIndex, tid); \
load_coll(c, channel->devCollectives+nextIndex, tid, &abortCount); \
} \
}
#else
@@ -98,61 +206,19 @@ __global__ void NCCL_KERN_NAME(coll, op, dtype)(struct ncclColl firstColl) { \
IMPL_COLL_KERN(coll##LL, op, ncclFunc, dtype, ctype, FUNC_INDEX(ncclColl, ncclOp, ncclType, 1, al)) \
#define IMPL_COLL3(coll, op, ncclFunc, dtype, ctype, ncclColl, ncclOp, ncclType) \
IMPL_COLL4(coll##Ring, op, ncclFunc, dtype, ctype, ncclColl, ncclOp, ncclType, 0) \
IMPL_COLL4(coll##Tree, op, ncclFunc, dtype, ctype, ncclColl, ncclOp, ncclType, 1)
IMPL_COLL4(coll##Ring, op, ncclFunc, dtype, ctype, ncclColl, ncclOp, ncclType, 0)
#if NCCL_TYPE == 0
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, i8, int8_t, ncclColl, ncclOp, ncclInt8)
#elif NCCL_TYPE == 1
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, u8, uint8_t, ncclColl, ncclOp, ncclUint8)
#elif NCCL_TYPE == 2
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, i32, int32_t, ncclColl, ncclOp, ncclInt32)
#elif NCCL_TYPE == 3
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, u32, uint32_t, ncclColl, ncclOp, ncclUint32)
#elif NCCL_TYPE == 4
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, i64, int64_t, ncclColl, ncclOp, ncclInt64)
#elif NCCL_TYPE == 5
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, u64, uint64_t, ncclColl, ncclOp, ncclUint64)
#elif NCCL_TYPE == 6
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, f16, half, ncclColl, ncclOp, ncclFloat16)
#elif NCCL_TYPE == 7
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, f32, float, ncclColl, ncclOp, ncclFloat32)
#elif NCCL_TYPE == 8
#define IMPL_COLL2(coll, op, ncclFunc, ncclColl, ncclOp) \
IMPL_COLL3(coll, op, ncclFunc, i8, int8_t, ncclColl, ncclOp, ncclInt8) \
IMPL_COLL3(coll, op, ncclFunc, u8, uint8_t, ncclColl, ncclOp, ncclUint8) \
IMPL_COLL3(coll, op, ncclFunc, i32, int32_t, ncclColl, ncclOp, ncclInt32) \
IMPL_COLL3(coll, op, ncclFunc, u32, uint32_t, ncclColl, ncclOp, ncclUint32) \
IMPL_COLL3(coll, op, ncclFunc, i64, int64_t, ncclColl, ncclOp, ncclInt64) \
IMPL_COLL3(coll, op, ncclFunc, u64, uint64_t, ncclColl, ncclOp, ncclUint64) \
IMPL_COLL3(coll, op, ncclFunc, f16, half, ncclColl, ncclOp, ncclFloat16) \
IMPL_COLL3(coll, op, ncclFunc, f32, float, ncclColl, ncclOp, ncclFloat32) \
IMPL_COLL3(coll, op, ncclFunc, f64, double, ncclColl, ncclOp, ncclFloat64)
#endif
// Reduction define all functions
#if NCCL_OP == 0
#define IMPL_COLL_R(collf, colln) \
IMPL_COLL2(collf, sum, FuncSum, colln, ncclSum);
#elif NCCL_OP == 1
#define IMPL_COLL_R(collf, colln) \
IMPL_COLL2(collf, prod, FuncProd, colln, ncclProd);
#elif NCCL_OP == 2
#define IMPL_COLL_R(collf, colln) \
IMPL_COLL2(collf, min, FuncMin, colln, ncclMin);
#elif NCCL_OP == 3
#define IMPL_COLL_R(collf, colln) \
IMPL_COLL2(collf, max, FuncMax, colln, ncclMax);
#endif
// Copy primitives only define one
#if NCCL_OP == 0 && NCCL_TYPE == 0
#define IMPL_COLL_C(collf, colln) \
IMPL_COLL3(collf, copy, FuncSum, i8, int8_t, colln, ncclSum, ncclInt8);
#else
#define IMPL_COLL_C(collf, colln)
#endif
#define COLL_UNROLL 4
#define COLL_UNROLL 2
#endif
src/collectives/device/common_kernel.h
+31 -6
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, 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,25 @@
#include <cstdio>
#include <cstdint>
#include <cuda_runtime.h>
#include <hip/hip_runtime.h>
// Define min for ssize_t
static __device__ int min(int a, ssize_t b) { return (a < b) ? a : b; }
typedef uint64_t PackType;
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
template<class FUNC, typename T>
struct MULTI {
__device__ PackType operator()(const PackType x, const PackType y) const
{
return FUNC()(x, y);
}
};
#else
// unpack x and y to elements of type T and apply FUNC to each element
template<class FUNC, typename T>
struct MULTI {
@@ -192,6 +205,8 @@ struct MULTI<FUNC, int64_t> {
}
};
#endif //defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
template<typename T> inline __device__
T vFetch(const volatile T* ptr) {
return *ptr;
@@ -202,7 +217,7 @@ void vStore(volatile T* ptr, const T val) {
*ptr = val;
}
#if CUDART_VERSION < 9000
#if CUDART_VERSION < 9000 && !(defined(__HIP_PLATFORM_HCC__) || defined(__HCC__))
template<> inline __device__
half vFetch<half>(const volatile half* ptr) {
half r;
@@ -239,14 +254,24 @@ struct MULTI128 {
};
inline __device__ void Fetch128(Pack128& v, const Pack128* p) {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
v.x = p->x;
v.y = p->y;
#else
asm volatile("ld.volatile.global.v2.u64 {%0,%1}, [%2];" : "=l"(v.x), "=l"(v.y) : "l"(p) : "memory");
#endif
}
inline __device__ void Store128(Pack128* p, Pack128& v) {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
p->x = v.x;
p->y = v.y;
#else
asm volatile("st.volatile.global.v2.u64 [%0], {%1,%2};" :: "l"(p), "l"(v.x), "l"(v.y) : "memory");
#endif
}
template<class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
__device__ __forceinline__ void ReduceCopyMulti(const int tid, const int nthreads,
__device__ void ReduceCopyMulti(const int tid, const int nthreads,
int nsrcs, const T* srcs[MAXSRCS], int ndsts, T* dsts[MAXDSTS],
const int offset, const int N) {
for (int idx = offset+tid; idx < offset+N; idx += nthreads) {
@@ -263,10 +288,10 @@ __device__ __forceinline__ void ReduceCopyMulti(const int tid, const int nthread
}
}
#define WARP_SIZE 32
#define WARP_SIZE 64
template<class FUNC, typename T, int UNROLL, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
__device__ __forceinline__ void ReduceCopy128bMulti( const int w, const int nw, const int t,
__device__ void ReduceCopy128bMulti( const int w, const int nw, const int t,
int nsrcs, const T* s[MAXSRCS], int ndsts, T* d[MAXDSTS],
const int elemOffset, const int Npack) {
const int inc = nw * UNROLL * WARP_SIZE;
@@ -316,7 +341,7 @@ __device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(Pack128); }
#define AUTOUNROLL (UNROLL*(4/(MINDSTS+MINSRCS)))
template<int UNROLL, class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
__device__ __forceinline__ void ReduceOrCopyMulti(const int tid, const int nthreads,
__device__ void ReduceOrCopyMulti(const int tid, const int nthreads,
int nsrcs, const T* srcs[MAXSRCS], int ndsts, T* dsts[MAXDSTS],
int N) {
int Nrem = N;
src/collectives/device/functions.cu
+1 -64
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -8,69 +9,5 @@
#include "collectives.h"
#include "common.h"
#define NCCL_FUNC5(coll, op, dtype) \
NCCL_COLL_NAME(coll, op, dtype), \
NCCL_COLL_NAME(coll##LL, op, dtype)
#define NCCL_FUNC4(coll, op, dtype) \
NCCL_FUNC5(coll##Ring, op, dtype), \
NCCL_FUNC5(coll##Tree, op, dtype)
// Must be consistent with ncclDataType_t
#define NCCL_FUNCS3A(coll, op) \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, u8), \
NCCL_FUNC4(coll, op, i32), \
NCCL_FUNC4(coll, op, u32), \
NCCL_FUNC4(coll, op, i64), \
NCCL_FUNC4(coll, op, u64), \
NCCL_FUNC4(coll, op, f16), \
NCCL_FUNC4(coll, op, f32), \
NCCL_FUNC4(coll, op, f64)
#define NCCL_FUNCS3B(coll, op) \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8), \
NCCL_FUNC4(coll, op, i8)
// Must be consistent with ncclRedOp_t
#define NCCL_FUNCS2A(coll) \
NCCL_FUNCS3A(coll, sum ), \
NCCL_FUNCS3A(coll, prod), \
NCCL_FUNCS3A(coll, max ), \
NCCL_FUNCS3A(coll, min )
#define NCCL_FUNCS2B(coll) \
NCCL_FUNCS3B(coll, copy), \
NCCL_FUNCS3B(coll, copy), \
NCCL_FUNCS3B(coll, copy), \
NCCL_FUNCS3B(coll, copy)
// Must be consistent with ncclColl_t
#define NCCL_FUNCS() { \
NCCL_FUNCS2B(ncclBroadcast), \
NCCL_FUNCS2A(ncclReduce), \
NCCL_FUNCS2B(ncclAllGather), \
NCCL_FUNCS2A(ncclReduceScatter), \
NCCL_FUNCS2A(ncclAllReduce) }
// Must be consistent with the ncclFuncSet enum
__device__ ncclKern_t ncclFuncs[ncclCollCount*ncclNumOps*ncclNumTypes*2*2] = {
// Don't try to initialize the host shadow copy of this device-side global
// variable. There is no host pointer to a device-side function, which
// confuses clang. This will be fixed in the next clang release.
#if __CUDA_ARCH__
NCCL_FUNCS2B(ncclBroadcast),
NCCL_FUNCS2A(ncclReduce),
NCCL_FUNCS2B(ncclAllGather),
NCCL_FUNCS2A(ncclReduceScatter),
NCCL_FUNCS2A(ncclAllReduce)
#endif
};
// Workaround for https://reviews.llvm.org/D55580
__device__ void ncclWorkaroundClangD55580() {}
src/collectives/device/primitives.h
+186 -113
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -51,24 +52,29 @@ class ncclPrimitives {
const T* recvBuff[NRECV];
T* sendBuff[NSEND];
struct ncclDevComm* comm;
uint32_t* abortCount;
inline __device__ int recvOffset(int i) { return (recvStep[i]%NCCL_STEPS)*stepSize; }
inline __device__ int sendOffset(int i) { return (sendStep[i]%NCCL_STEPS)*stepSize; }
inline __device__ const T* recvPtr(int i) { return ((const T*)recvBuff[i])+recvOffset(i); }
inline __device__ T* sendPtr(int i) { return ((T*)sendBuff[i])+sendOffset(i); }
__device__ int recvOffset(int i) { return (recvStep[i]%NCCL_STEPS)*stepSize; }
__device__ int sendOffset(int i) { return (sendStep[i]%NCCL_STEPS)*stepSize; }
__device__ const T* recvPtr(int i) { return ((const T*)recvBuff[i])+recvOffset(i); }
__device__ T* sendPtr(int i) { return ((T*)sendBuff[i])+sendOffset(i); }
inline __device__ void barrier() {
__device__ void barrier() {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
__syncthreads();
#else
asm volatile ("bar.sync 1, %0;" :: "r"(nthreads));
#endif
}
uint32_t mismatch = 0;
const uint64_t opCount;
inline __device__ void checkMismatch(volatile uint64_t* remoteOpCount) {
__device__ void checkMismatch(volatile uint64_t* remoteOpCount) {
if (mismatch) {
// In non-LL, we use _threadfence_system before incrementing opCount, yet we are still waiting for credits here, so there must be a size mismatch
*(comm->fatalDevError) = ncclDevAssertedMismatch;
} else if (remoteOpCount && *remoteOpCount > opCount) {
STORE(comm->fatalDevError, ncclDevAssertedMismatch);
} else if (remoteOpCount && LOAD(remoteOpCount) > opCount) {
mismatch += 1;
}
}
@@ -76,63 +82,78 @@ class ncclPrimitives {
uint32_t spins = 0;
uint32_t abort = 0;
inline __device__ int checkAbort(volatile uint64_t* remoteOpCount) {
__device__ int checkAbort(volatile uint64_t* remoteOpCount) {
spins++;
abort = LOAD(comm->abortFlag);
if (spins == SPINS_BEFORE_CHECK_ABORT) {
abort = *(comm->abortFlag);
checkMismatch(remoteOpCount);
spins = 0;
}
return abort;
}
inline __device__ void waitRecv(int i) {
__device__ void waitRecv(int i) {
spins = 0;
mismatch = 0;
recvStep[i] += SLICESTEPS;
if (tid == i) {
while (*(waitPtr) < recvStep[i]) {
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t t0 = clock64();
#endif
while (LOAD(waitPtr) < recvStep[i]) {
if (checkAbort(recvConn[i]->opCountRem)) break;
}
#ifdef ENABLE_PROFILING
__atomic_fetch_add(&devProf->wait_recv_cycle[blockIdx.x], clock64() - t0, __ATOMIC_SEQ_CST);
#endif
}
}
inline __device__ void waitSend(int i) {
__device__ void waitSend(int i) {
spins = 0;
mismatch = 0;
sendStep[i] += SLICESTEPS;
if (tid == WARP_SIZE+i) {
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t t0 = clock64();
#endif
while (sendConnHead[i] + NCCL_STEPS < sendStep[i]) {
sendConnHead[i] = *waitPtr;
sendConnHead[i] = LOAD(waitPtr);
if (checkAbort(sendConn[i]->opCountRem)) break;
}
#ifdef ENABLE_PROFILING
__atomic_fetch_add(&devProf->wait_send_cycle[blockIdx.x], clock64() - t0, __ATOMIC_SEQ_CST);
#endif
}
}
inline __device__ void postRecv(int i) {
*(recvConn[i]->head) = recvStep[i] += SLICESTEPS;
STORE(recvConn[i]->head, recvStep[i]);
}
inline __device__ void postSend(int i) {
*(sendConn[i]->tail) = sendStep[i] += SLICESTEPS;
if (sendConn[i]->next_hdp_reg) STORE(sendConn[i]->next_hdp_reg, 0x1);
STORE(sendConn[i]->tail, sendStep[i]);
}
inline __device__ void postSendSize(int i, int size) {
if (sendConn[i]->fifo) sendConn[i]->fifo[sendStep[i]%NCCL_STEPS] = size;
__device__ void postSendSize(int i, int size) {
if (sendConn[i]->fifo) STORE(sendConn[i]->fifo+((sendStep[i]-SLICESTEPS)%NCCL_STEPS), size);
}
template <int DIRECTRECV>
inline __device__ const T* directRecvPtr(int i, int directOffset) {
__device__ const T* directRecvPtr(int i, int directOffset) {
return DIRECTRECV && recvDirectBuff[i] ? recvDirectBuff[i]+directOffset : recvPtr(i);
}
template <int DIRECTSEND>
inline __device__ T* directSendPtr(int i, int directOffset) {
__device__ T* directSendPtr(int i, int directOffset) {
return DIRECTSEND && sendDirectBuff[i] ? sendDirectBuff[i]+directOffset : sendPtr(i);
}
template <int DIRECTRECV, int DIRECTSEND, int RECV, int SEND, int SRC, int DST>
inline __device__ void
__device__ void
GenericOp(const T* srcPtr, T* dstPtr, int nelem, int directOffset) {
int offset = 0;
int sliceSize = stepSize * SLICESTEPS;
@@ -154,157 +175,155 @@ class ncclPrimitives {
#pragma unroll 1
for (int slice=0; slice<SLICESPERCHUNK; ++slice) {
int realSize = max(0, min(sliceSize, nelem-offset));
if (tid < nthreads) {
FOR_SEND(waitSend);
FOR_RECV(waitRecv);
if (realSize > 0) {
barrier();
if (DIRECTRECV && recvDirectBuff[0]) {
// We can only have one direct receive. Since srcs[0] == dstPtr+offset, skip one copy
if (SEND) {
ReduceOrCopyMulti<UNROLL, FUNC, T, 1, 1, 1, NSEND>(tid, nthreads, 1, srcs, nsend, dsts+1, realSize);
}
} else {
ReduceOrCopyMulti<UNROLL, FUNC, T, RECV+SRC, RECV*NRECV+SRC, SEND+DST, SEND*NSEND+DST>(tid, nthreads, RECV*nrecv+SRC, srcs, SEND*nsend+DST, dsts, realSize);
FOR_SEND(waitSend);
FOR_RECV(waitRecv);
if (realSize > 0) {
barrier();
if (DIRECTRECV && recvDirectBuff[0]) {
// We can only have one direct receive. Since srcs[0] == dstPtr+offset, skip one copy
if (SEND) {
ReduceOrCopyMulti<UNROLL, FUNC, T, 1, 1, 1, NSEND>(tid, nthreads, 1, srcs, nsend, dsts+1, realSize);
}
} else {
ReduceOrCopyMulti<UNROLL, FUNC, T, RECV+SRC, RECV*NRECV+SRC, SEND+DST, SEND*NSEND+DST>(tid, nthreads, RECV*nrecv+SRC, srcs, SEND*nsend+DST, dsts, realSize);
}
exitIfAbortBarrier(abort);
} else {
exitIfAbortBarrier(abort);
FOR_SEND(postSendSize, realSize*sizeof(T));
if (SEND) __threadfence_system();
FOR_SEND(postSend);
FOR_RECV(postRecv);
}
for (int i=0; i<RECV*NRECV+SRC; i++) srcs[i] += sliceSize;
for (int i=0; i<SEND*NSEND+DST; i++) dsts[i] += sliceSize;
offset += sliceSize;
exitIfAbortBarrier(abort, abortCount);
if (tid == 0) FOR_SEND(postSendSize, realSize*sizeof(T));
if (SEND) __threadfence_system();
if (tid == 0) FOR_SEND(postSend);
if (tid == 0) FOR_RECV(postRecv);
}
for (int i=0; i<RECV*NRECV+SRC; i++) srcs[i] += sliceSize;
for (int i=0; i<SEND*NSEND+DST; i++) dsts[i] += sliceSize;
offset += sliceSize;
}
__device__ __forceinline__ void loadRecvConn(struct ncclConnInfo* conn, int i, T* directBuff) {
__device__ void loadRecvConn(struct ncclConnInfo* conn, int i, T* directBuff) {
recvConn[i] = conn;
recvBuff[i] = (const T*)recvConn[i]->buff;
recvStep[i] = recvConn[i]->step;
recvBuff[i] = (const T*)LOAD(&recvConn[i]->buff);
recvStep[i] = LOAD(&recvConn[i]->step);
recvStep[i] = ROUNDUP(recvStep[i], SLICESPERCHUNK*SLICESTEPS);
// Return credits in case we rounded up.
if (tid == nthreads) *recvConn[i]->head = recvStep[i];
if (tid == 0) STORE(recvConn[i]->head, recvStep[i]);
if (tid == i) {
waitPtr = recvConn[i]->tail;
*(recvConn[i]->opCountLoc) = opCount;
waitPtr = LOAD(&recvConn[i]->tail);
STORE(recvConn[i]->opCountLoc, opCount);
}
recvDirectBuff[i] = NULL;
if (directBuff && recvConn[i]->direct) {
recvDirectBuff[i] = directBuff;
if (tid == 0) *recvConn[i]->ptrExchange = directBuff;
if (tid == 0) STORE(recvConn[i]->ptrExchange, directBuff);
}
nrecv++;
}
__device__ __forceinline__ void loadSendConn(struct ncclConnInfo* conn, int i, T* directBuff) {
__device__ void loadSendConn(struct ncclConnInfo* conn, int i, T* directBuff) {
sendConn[i] = conn;
sendBuff[i] = (T*)sendConn[i]->buff;
sendStep[i] = sendConn[i]->step;
sendBuff[i] = (T*)LOAD(&sendConn[i]->buff);
sendStep[i] = LOAD(&sendConn[i]->step);
sendStep[i] = ROUNDUP(sendStep[i], SLICESPERCHUNK*SLICESTEPS);
if (tid == WARP_SIZE+i) {
waitPtr = sendConn[i]->head;
sendConnHead[i] = *waitPtr;
*(sendConn[i]->opCountLoc) = opCount;
waitPtr = LOAD(&sendConn[i]->head);
sendConnHead[i] = LOAD(waitPtr);
STORE(sendConn[i]->opCountLoc, opCount);
}
sendDirectBuff[i] = NULL;
if (directBuff && sendConn[i]->direct) {
void* volatile* ptr = sendConn[i]->ptrExchange;
while ((sendDirectBuff[i] = (T*)(*ptr)) == NULL);
while ((sendDirectBuff[i] = (T*)(LOAD(ptr))) == NULL);
__syncthreads();
if (tid == 0) *ptr = NULL;
if (tid == 0) STORE(ptr, NULL);
}
nsend++;
}
__device__ __forceinline__ void saveRecvConn(int i) {
__device__ void saveRecvConn(int i) {
if (tid == i) {
recvConn[i]->step = recvStep[i];
STORE(&recvConn[i]->step, recvStep[i]);
__threadfence_system();
*(recvConn[i]->opCountLoc) += 1;
__atomic_fetch_add(recvConn[i]->opCountLoc, 1, __ATOMIC_SEQ_CST);
}
}
__device__ __forceinline__ void saveSendConn(int i) {
__device__ void saveSendConn(int i) {
if (tid == WARP_SIZE+i) {
sendConn[i]->step = sendStep[i];
STORE(&sendConn[i]->step, sendStep[i]);
__threadfence_system();
*(sendConn[i]->opCountLoc) += 1;
__atomic_fetch_add(sendConn[i]->opCountLoc, 1, __ATOMIC_SEQ_CST);
}
}
public:
__device__ __forceinline__
__device__
ncclPrimitives(const int tid, const int nthreads, int* recvPeers, int* sendPeers, T* directBuff, int stepSize, struct ncclChannel* channel, struct ncclDevComm* comm, const uint64_t opCount)
: comm(comm), tid(tid), nthreads(nthreads), stepSize(stepSize), opCount(opCount) {
// Make sure step is updated before we read it
abortCount = channel->abortCount;
__syncthreads();
for (int i=0; i<NRECV && recvPeers[i] >= 0; i++) loadRecvConn(&channel->devPeers[recvPeers[i]].recv.conn, i, directBuff);
for (int i=0; i<NSEND && sendPeers[i] >= 0; i++) loadSendConn(&channel->devPeers[sendPeers[i]].send.conn, i, directBuff);
// disable directBuff
for (int i=0; i<NRECV && recvPeers[i] >= 0; i++) loadRecvConn(&channel->devPeers[recvPeers[i]].recv.conn, i, 0);
for (int i=0; i<NSEND && sendPeers[i] >= 0; i++) loadSendConn(&channel->devPeers[sendPeers[i]].send.conn, i, 0);
}
__device__ __forceinline__ void
__device__ void
send(const T* src, int nelem) {
GenericOp<0, 0, 0, 1, 1, 0>(src, NULL, nelem, 0);
}
__device__ __forceinline__ void
__device__ void
directSend(const T* src, int directOffset, int nelem) {
GenericOp<0, 1, 0, 1, 1, 0>(src, NULL, nelem, directOffset);
}
__device__ __forceinline__ void
__device__ void
recv(T* dst, int nelem) {
GenericOp<0, 0, 1, 0, 0, 1>(NULL, dst, nelem, 0);
}
__device__ __forceinline__ void
__device__ void
directRecv(T* dst, int directOffset, int nelem) {
GenericOp<1, 0, 1, 0, 0, 1>(NULL, dst, nelem, directOffset);
}
__device__ __forceinline__ void
__device__ void
copySend(const T* src, T* dst, int nelem) {
GenericOp<0, 0, 0, 1, 1, 1>(src, dst, nelem, 0);
}
__device__ __forceinline__ void
__device__ void
directCopySend(const T* src, T* dst, int directOffset, int nelem) {
GenericOp<0, 1, 0, 1, 1, 1>(src, dst, nelem, directOffset);
}
__device__ __forceinline__ void
__device__ void
recvCopySend(T* dst, int nelem) {
GenericOp<0, 0, 1, 1, 0, 1>(NULL, dst, nelem, 0);
}
__device__ __forceinline__ void
__device__ void
directRecvCopySend(T* dst, int directOffset, int nelem) {
GenericOp<1, 1, 1, 1, 0, 1>(NULL, dst, nelem, directOffset);
}
__device__ __forceinline__ void
__device__ void
recvReduceCopy(const T* src, T* dst, int nelem) {
GenericOp<0, 0, 1, 0, 1, 1>(src, dst, nelem, 0);
}
__device__ __forceinline__ void
__device__ void
recvReduceSend(const T* src, int nelem) {
GenericOp<0, 0, 1, 1, 1, 0>(src, NULL, nelem, 0);
}
__device__ __forceinline__ void
__device__ void
recvReduceCopySend(const T* src, T* dst, int nelem) {
GenericOp<0, 0, 1, 1, 1, 1>(src, dst, nelem, 0);
}
__device__ __forceinline__ void
__device__ void
directRecvReduceCopySend(const T* src, T* dst, int directOffset, int nelem) {
// Direct is only for the send part
GenericOp<0, 1, 1, 1, 1, 1>(src, dst, nelem, directOffset);
}
__device__ __forceinline__ ~ncclPrimitives() {
__device__ ~ncclPrimitives() {
// Save steps for next collective. Have thread 0 do it to be compatible
// with the way LL works.
for (int i=0; i<NRECV && i<nrecv; i++) saveRecvConn(i);
@@ -330,19 +349,22 @@ class ncclLLPrimitives {
union ncclLLFifoLine* recvBuff[NRECV];
union ncclLLFifoLine* sendBuff[NSEND];
struct ncclDevComm* comm;
uint32_t* abortCount;
inline __device__ int recvOffset(int i) { return (recvStep[i]%NCCL_STEPS)*NCCL_LL_SLICE_LINES; }
inline __device__ int sendOffset(int i) { return (sendStep[i]%NCCL_STEPS)*NCCL_LL_SLICE_LINES; }
inline __device__ union ncclLLFifoLine* recvPtr(int i) { return recvBuff[i]+recvOffset(i); }
inline __device__ union ncclLLFifoLine* sendPtr(int i) { return sendBuff[i]+sendOffset(i); }
inline __device__ uint32_t recvFlag(int i) { return NCCL_LL_FLAG(recvStep[i]+1); }
inline __device__ uint32_t sendFlag(int i) { return NCCL_LL_FLAG(sendStep[i]+1); }
__device__ int recvOffset(int i) { return (recvStep[i]%NCCL_STEPS)*NCCL_LL_SLICE_LINES; }
__device__ int sendOffset(int i) { return (sendStep[i]%NCCL_STEPS)*NCCL_LL_SLICE_LINES; }
__device__ union ncclLLFifoLine* recvPtr(int i) { return recvBuff[i]+recvOffset(i); }
__device__ union ncclLLFifoLine* sendPtr(int i) { return sendBuff[i]+sendOffset(i); }
__device__ uint32_t recvFlag(int i) { return NCCL_LL_FLAG(recvStep[i]+1); }
__device__ uint32_t sendFlag(int i) { return NCCL_LL_FLAG(sendStep[i]+1); }
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#else
// Exit If Abort Barrier : make sure all threads exit consistently
// Each thread sets a predicate to true if val == 1
// all CTA's threads enter the barrier and do a popc on their predicates being True
// If any of the thread's predicate was True, all the threads call exit()
inline __device__ void exitIfAbortLocalBarrier() {
__device__ void exitIfAbortLocalBarrier() {
uint32_t popc;
asm ("{");
asm volatile (" .reg .pred barr_pred;");
@@ -354,20 +376,25 @@ class ncclLLPrimitives {
exitIfAbortBarrier(1);
}
}
#endif
inline __device__ void barrier() {
__device__ void barrier() {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
__syncthreads();
#else
asm volatile ("bar.sync 1, %0;" :: "r"(nthreads));
#endif
}
uint32_t mismatch = 0;
const uint64_t opCount;
inline __device__ void checkMismatch(volatile uint64_t* remoteOpCount) {
__device__ void checkMismatch(volatile uint64_t* remoteOpCount) {
if (mismatch > 20) {
// We have seen that the peer advanced opcount so many times yet we are still waiting for credit of current op, so it is _most likely_ a mismatch
// Note that we are not using _threadfence_system in LL so the error cannot be asserted
*(comm->fatalDevError) = ncclDevSuspectedMismatch;
} else if (remoteOpCount && *remoteOpCount > opCount) {
STORE(comm->fatalDevError, ncclDevSuspectedMismatch);
} else if (remoteOpCount && LOAD(remoteOpCount) > opCount) {
mismatch += 1;
}
}
@@ -375,37 +402,37 @@ class ncclLLPrimitives {
uint32_t spins = 0;
uint32_t abort = 0;
inline __device__ int checkAbort(volatile uint64_t* remoteOpCount) {
__device__ int checkAbort(volatile uint64_t* remoteOpCount) {
spins++;
abort = LOAD(comm->abortFlag);
if (spins == SPINS_BEFORE_CHECK_ABORT) {
abort = *(comm->abortFlag);
checkMismatch(remoteOpCount);
spins = 0;
}
return abort;
}
inline __device__ void waitSend(int i, int nbytes) {
__device__ void waitSend(int i, int nbytes) {
spins = 0;
mismatch = 0;
if (tid == WARP_SIZE+i) {
while (sendConnHead + NCCL_STEPS < sendStep[i] + 1) {
sendConnHead = *waitPtr;
sendConnHead = LOAD(waitPtr);
if (checkAbort(sendConn[i]->opCountRem)) break;
}
if (fifoPtr) {
int size = ((sendStep[i] & NCCL_LL_CLEAN_MASK) == NCCL_LL_CLEAN_MASK) ? NCCL_LL_SLICE_LINES*sizeof(union ncclLLFifoLine) : nbytes;
fifoPtr[sendStep[i]%NCCL_STEPS] = size;
STORE(fifoPtr+sendStep[i]%NCCL_STEPS, size);
}
}
}
inline __device__ void postRecv(int i) {
__device__ void postRecv(int i) {
recvStep[i]++;
if (tid == i) *postPtr = recvStep[i];
if (tid == i) STORE(postPtr, recvStep[i]);
}
inline __device__ void postSend(int i, int offset) {
__device__ void postSend(int i, int offset) {
// LL Cleanup : write all flags in the slice to make sure we don't have
// data corruption when flag loops over.
if ((sendStep[i] & NCCL_LL_CLEAN_MASK) == NCCL_LL_CLEAN_MASK) {
@@ -414,22 +441,46 @@ class ncclLLPrimitives {
sendStep[i]++;
}
__device__ uint64_t readLL(int i, int offset) {
__device__ __attribute__((noinline)) uint64_t readLL(int i, int offset) {
union ncclLLFifoLine* src = recvPtr(i) + offset;
uint32_t flag = recvFlag(i);
uint32_t data1, flag1, data2, flag2;
spins = 0;
mismatch = 0;
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
using Vec = uint32_t __attribute__((ext_vector_type(4)));
Vec i4;
do {
asm volatile ("flat_load_dwordx4 %0, %1, glc\n"
"s_waitcnt vmcnt(0)\n"
"buffer_wbinvl1_vol\n" : "=v"(i4) : "v"(src));
if (checkAbort(recvConn[i]->opCountRem)) break;
} while (i4[1] != flag || i4[3] != flag);
uint64_t val64 = (uint64_t)(i4[0]) + (((uint64_t)i4[2]) << 32);
#else
do {
asm volatile("ld.volatile.global.v4.u32 {%0,%1,%2,%3}, [%4];" : "=r"(data1), "=r"(flag1), "=r"(data2), "=r"(flag2) : "l"(&src->i4));
if (checkAbort(recvConn[i]->opCountRem)) break;
} while ((flag1 != flag) || (flag2 != flag));
uint64_t val64 = data1 + (((uint64_t)data2) << 32);
#endif
return val64;
}
__device__ void storeLL(union ncclLLFifoLine* dst, uint64_t val, uint32_t flag) {
__device__ __attribute__((noinline)) void storeLL(union ncclLLFifoLine* dst, uint64_t val, uint32_t flag) {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
using Vec = uint32_t __attribute__((ext_vector_type(4)));
Vec i4;
i4[0] = val & 0xffffffff;
i4[1] = flag;
i4[2] = (val >> 32);
i4[3] = flag;
asm volatile ("flat_store_dwordx4 %0, %1, glc\n"
"s_waitcnt vmcnt(0)\n"
"buffer_wbinvl1_vol\n" : : "v"(dst), "v"(i4));
#else
asm volatile("st.volatile.global.v4.u32 [%0], {%1,%2,%3,%4};" :: "l"(&dst->i4), "r"((uint32_t)val), "r"(flag), "r"((uint32_t)(val >> 32)), "r"(flag));
#endif
}
// Using memcpy handles misaligned pointers.
@@ -453,7 +504,7 @@ class ncclLLPrimitives {
uint64_t* dstPack = (uint64_t*)dstPtr;
int offset = tid;
// Do multiples of 64 bits
#pragma unroll 2
#pragma unroll 1
for (; offset<npack; offset+=nthreads) {
// Recv : local, then intra-node, then inter-node
uint64_t val = SRC ? readAL(srcPack+offset) : readLL(0, offset);
@@ -478,56 +529,61 @@ class ncclLLPrimitives {
}
}
}
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
exitIfAbortBarrier(abort, abortCount);
#else
exitIfAbortLocalBarrier();
#endif
FOR_RECV(postRecv);
FOR_SEND(postSend, offset);
}
__device__ __forceinline__ void loadRecvConn(struct ncclConnInfo* conn, int i) {
__device__ void loadRecvConn(struct ncclConnInfo* conn, int i) {
recvConn[i] = conn;
recvBuff[i] = recvConn[i]->llBuff;
recvStep[i] = recvConn[i]->step;
if (tid == i) {
postPtr = recvConn[i]->head;
*(recvConn[i]->opCountLoc) = opCount;
STORE(recvConn[i]->opCountLoc, opCount);
}
nrecv++;
}
__device__ __forceinline__ void loadSendConn(struct ncclConnInfo* conn, int i) {
__device__ void loadSendConn(struct ncclConnInfo* conn, int i) {
sendConn[i] = conn;
sendBuff[i] = sendConn[i]->llBuff;
sendStep[i] = sendConn[i]->step;
if (tid == WARP_SIZE+i) {
waitPtr = sendConn[i]->head;
fifoPtr = sendConn[i]->fifo;
sendConnHead = *waitPtr;
*(sendConn[i]->opCountLoc) = opCount;
sendConnHead = LOAD(waitPtr);
STORE(sendConn[i]->opCountLoc, opCount);
}
nsend++;
}
__device__ __forceinline__ void saveRecvConn(int i) {
__device__ void saveRecvConn(int i) {
if (tid == i) {
recvConn[i]->step = recvStep[i];
*(recvConn[i]->opCountLoc) += 1;
__atomic_fetch_add(recvConn[i]->opCountLoc, 1, __ATOMIC_SEQ_CST);
__threadfence_block();
}
}
__device__ __forceinline__ void saveSendConn(int i) {
__device__ void saveSendConn(int i) {
if (tid == WARP_SIZE+i) {
sendConn[i]->step = sendStep[i];
*(sendConn[i]->opCountLoc) += 1;
__atomic_fetch_add(sendConn[i]->opCountLoc, 1, __ATOMIC_SEQ_CST);
__threadfence_block();
}
}
public:
__device__ __forceinline__
__device__
ncclLLPrimitives(const int tid, const int nthreads, int* recvPeers, int* sendPeers, struct ncclChannel* channel, struct ncclDevComm* comm, const uint64_t opCount)
: comm(comm), tid(tid), nthreads(nthreads), opCount(opCount) {
// Make sure step is updated before we read it.
abortCount = channel->abortCount;
barrier();
for (int i=0; i<NRECV && recvPeers[i] >= 0; i++) loadRecvConn(&channel->devPeers[recvPeers[i]].recv.conn, i);
@@ -562,10 +618,27 @@ class ncclLLPrimitives {
return LLGenericOp<1, 1, 1, 1>(src, dst, nelem);
}
__device__ __forceinline__ ~ncclLLPrimitives() {
__device__ ~ncclLLPrimitives() {
// Save steps for the next operation
for (int i=0; i<NRECV && i<nrecv; i++) saveRecvConn(i);
for (int i=0; i<NSEND && i<nsend; i++) saveSendConn(i);
}
};
#ifdef ENABLE_PROFILING
#define INIT_COUNTER \
if (tid==0) { t0 = clock64(); ws = LOAD(&(devProf->wait_send_cycle[blockIdx.x])); \
wr = LOAD(&(devProf->wait_recv_cycle[blockIdx.x])); }
#define ACCUMULATE_COUNTER(prim) \
if (tid==0) { __atomic_fetch_add(&(devProf->prim##_cycle), clock64() - t0 \
+ ws - LOAD(&(devProf->wait_send_cycle[blockIdx.x])) \
+ wr - LOAD(&(devProf->wait_recv_cycle[blockIdx.x])), \
__ATOMIC_SEQ_CST); \
__atomic_fetch_add(&(devProf->prim##_byte), nelem * sizeof(T), __ATOMIC_SEQ_CST); }
#else
#define INIT_COUNTER
#define ACCUMULATE_COUNTER(prim)
#endif
#endif
src/collectives/device/reduce.cu
+9 -3
View File
@@ -1,11 +1,17 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* 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
************************************************************************/
#include "reduce.h"
#include "common.h"
#include "reduce.h"
#include "collectives.h"
IMPL_COLL_R(ncclReduce, ncclCollReduce);
#define UNROLL 4
IMPL_COLL2(ncclReduce, sum, FuncSum, ncclCollReduce, ncclSum);
IMPL_COLL2(ncclReduce, prod, FuncProd, ncclCollReduce, ncclProd);
IMPL_COLL2(ncclReduce, min, FuncMin, ncclCollReduce, ncclMin);
IMPL_COLL2(ncclReduce, max, FuncMax, ncclCollReduce, ncclMax);
src/collectives/device/reduce.h
+6 -1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -9,9 +10,10 @@
#include "collectives.h"
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceRingKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int nthreads = blockDim.x - 1;
const int nthreads = blockDim.x;
const int bid = args->bid;
struct ncclDevComm* comm = args->comm;
struct ncclChannel* channel = comm->channels+blockIdx.x;
@@ -48,9 +50,11 @@ __device__ void ncclReduceRingKernel(struct CollectiveArgs* args) {
}
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceTreeKernel(struct CollectiveArgs* args) { }
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceRingLLKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int bid = args->bid;
@@ -92,4 +96,5 @@ __device__ void ncclReduceRingLLKernel(struct CollectiveArgs* args) {
}
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceTreeLLKernel(struct CollectiveArgs* args) { }
src/collectives/device/reduce_kernel.h
+139
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -18,6 +19,123 @@ struct FuncNull {
}
};
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
//we really don't need any specializations and we don't need
//to break things into uint32_t
template<typename T>
__device__ inline T ncclMinFunc(T x, T y) { return y < x ? y : x; }
template<typename T>
__device__ inline T ncclMaxFunc(T x, T y) { return y < x ? x : y; }
template<typename T>
class FuncBase {
protected:
static constexpr auto n = sizeof(PackType) / sizeof(T);
union Cvt {
using Vec = T __attribute__((ext_vector_type(n)));
PackType data;
Vec vec;
static_assert(sizeof(Vec) == sizeof(data), "Vec must be the same size of data.");
};
};
template<>
class FuncBase<half> {
protected:
static constexpr auto n = sizeof(PackType) / sizeof(_Float16);
union Cvt {
using Vec = _Float16 __attribute__((ext_vector_type(n)));
PackType data;
Vec vec;
static_assert(sizeof(Vec) == sizeof(data), "Vec must be the same size of data.");
};
};
template<typename T>
struct FuncSum : private FuncBase<T> {
__device__ PackType operator()(PackType x, PackType y) const
{
using Cvt = typename FuncBase<T>::Cvt;
Cvt tmp_x{x};
tmp_x.vec += Cvt{y}.vec;
return tmp_x.data;
}
template<typename U = T, typename std::enable_if<!std::is_same<T, U>{}>* = nullptr>
__device__ T operator()(const T x, const T y) const {
return x + y;
}
};
template<typename T>
struct FuncProd : private FuncBase<T> {
__device__ PackType operator()(PackType x, PackType y) const
{
using Cvt = typename FuncBase<T>::Cvt;
Cvt tmp_x{x};
tmp_x.vec *= Cvt{y}.vec;
return tmp_x.data;
}
template<typename U = T, typename std::enable_if<!std::is_same<T, U>{}>* = nullptr>
__device__ T operator()(const T x, const T y) const {
return x * y;
}
};
template<typename T>
struct FuncMax : private FuncBase<T> {
__device__ PackType operator()(PackType x, PackType y) const
{
using Cvt = typename FuncBase<T>::Cvt;
Cvt tmp_x{x};
Cvt tmp_y{y};
for (auto i = 0u; i != FuncBase<T>::n; ++i) {
tmp_x.vec[i] = ncclMaxFunc(tmp_x.vec[i], tmp_y.vec[i]);
}
return tmp_x.data;
}
template<typename U = T, typename std::enable_if<!std::is_same<T, U>{}>* = nullptr>
__device__ T operator()(const T x, const T y) const {
return (x < y) ? y : x;
}
};
template<typename T>
struct FuncMin : private FuncBase<T> {
__device__ PackType operator()(PackType x, PackType y) const
{
using Cvt = typename FuncBase<T>::Cvt;
Cvt tmp_x{x};
Cvt tmp_y{y};
for (auto i = 0u; i != FuncBase<T>::n; ++i) {
tmp_x.vec[i] = ncclMinFunc(tmp_x.vec[i], tmp_y.vec[i]);
}
return tmp_x.data;
}
template<typename U = T, typename std::enable_if<!std::is_same<T, U>{}>* = nullptr>
__device__ T operator()(const T x, const T y) const {
return (x < y) ? x : y;
}
};
#else
template<typename T>
struct FuncSum {
__device__ T operator()(const T x, const T y) const {
@@ -62,12 +180,15 @@ static __device__ uint32_t addChar4(const uint32_t x, const uint32_t y) {
template<>
struct FuncSum<int8_t> {
__device__ uint32_t operator()(const uint32_t x, const uint32_t y) const {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#else
#if (__CUDA_ARCH__ >= 300) && (__CUDA_ARCH__ < 500)
int32_t rv, z=0;
asm("vadd4.s32.s32.s32 %0, %1, %2, %3;" : "=r"(rv) : "r"(x), "r"(y), "r"(z));
return rv;
#else
return addChar4(x, y);
#endif
#endif
}
__device__ int8_t operator()(const int8_t x, const int8_t y) const {
@@ -77,12 +198,15 @@ struct FuncSum<int8_t> {
template<>
struct FuncSum<uint8_t> {
__device__ uint32_t operator()(const uint32_t x, const uint32_t y) const {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#else
#if (__CUDA_ARCH__ >= 300) && (__CUDA_ARCH__ < 500)
int32_t rv, z=0;
asm("vadd4.u32.u32.u32 %0, %1, %2, %3;" : "=r"(rv) : "r"(x), "r"(y), "r"(z));
return rv;
#else
return addChar4(x, y);
#endif
#endif
}
__device__ uint8_t operator()(const uint8_t x, const uint8_t y) const {
@@ -126,6 +250,8 @@ template<>
struct FuncMax<int8_t> {
union converter { uint32_t storage; char4 a; };
__device__ uint32_t operator()(const uint32_t x, const uint32_t y) const {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#else
#if (__CUDA_ARCH__ >= 300) && (__CUDA_ARCH__ < 500)
int32_t rv, z=0;
asm("vmax4.s32.s32.s32 %0, %1, %2, %3;" : "=r"(rv) : "r"(x), "r"(y), "r"(z));
@@ -139,6 +265,7 @@ struct FuncMax<int8_t> {
cr.a.z = max(cx.a.z, cy.a.z);
cr.a.w = max(cx.a.w, cy.a.w);
return cr.storage;
#endif
#endif
}
__device__ int8_t operator()(const int8_t x, const int8_t y) const {
@@ -149,6 +276,8 @@ template<>
struct FuncMax<uint8_t> {
union converter { uint32_t storage; uchar4 a; };
__device__ uint32_t operator()(const uint32_t x, const uint32_t y) const {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#else
#if (__CUDA_ARCH__ >= 300) && (__CUDA_ARCH__ < 500)
int32_t rv, z=0;
asm("vmax4.u32.u32.u32 %0, %1, %2, %3;" : "=r"(rv) : "r"(x), "r"(y), "r"(z));
@@ -162,6 +291,7 @@ struct FuncMax<uint8_t> {
cr.a.z = max(cx.a.z, cy.a.z);
cr.a.w = max(cx.a.w, cy.a.w);
return cr.storage;
#endif
#endif
}
__device__ uint8_t operator()(const uint8_t x, const uint8_t y) const {
@@ -173,6 +303,8 @@ template<>
struct FuncMin<int8_t> {
union converter { uint32_t storage; char4 a; };
__device__ uint32_t operator()(const uint32_t x, const uint32_t y) const {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#else
#if (__CUDA_ARCH__ >= 300) && (__CUDA_ARCH__ < 500)
int32_t rv, z=0;
asm("vmin4.s32.s32.s32 %0, %1, %2, %3;" : "=r"(rv) : "r"(x), "r"(y), "r"(z));
@@ -186,6 +318,7 @@ struct FuncMin<int8_t> {
cr.a.z = min(cx.a.z, cy.a.z);
cr.a.w = min(cx.a.w, cy.a.w);
return cr.storage;
#endif
#endif
}
__device__ int8_t operator()(const int8_t x, const int8_t y) const {
@@ -196,6 +329,8 @@ template<>
struct FuncMin<uint8_t> {
union converter { uint32_t storage; uchar4 a; };
__device__ uint32_t operator()(const uint32_t x, const uint32_t y) const {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#else
#if (__CUDA_ARCH__ >= 300) && (__CUDA_ARCH__ < 500)
int32_t rv, z=0;
asm("vmin4.u32.u32.u32 %0, %1, %2, %3;" : "=r"(rv) : "r"(x), "r"(y), "r"(z));
@@ -209,6 +344,7 @@ struct FuncMin<uint8_t> {
cr.a.z = min(cx.a.z, cy.a.z);
cr.a.w = min(cx.a.w, cy.a.w);
return cr.storage;
#endif
#endif
}
__device__ uint8_t operator()(const uint8_t x, const uint8_t y) const {
@@ -299,4 +435,7 @@ struct FuncMin<half> {
return __float2half(fm);
}
};
#endif // defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)
#endif // REDUCE_KERNEL_H_
src/collectives/device/reduce_scatter.cu
+10 -3
View File
@@ -1,11 +1,18 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2015-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#include "reduce_scatter.h"
#include "common.h"
#include "reduce_scatter.h"
#include "collectives.h"
IMPL_COLL_R(ncclReduceScatter, ncclCollReduceScatter);
#define UNROLL 4
IMPL_COLL2(ncclReduceScatter, sum, FuncSum, ncclCollReduceScatter, ncclSum);
IMPL_COLL2(ncclReduceScatter, prod, FuncProd, ncclCollReduceScatter, ncclProd);
IMPL_COLL2(ncclReduceScatter, min, FuncMin, ncclCollReduceScatter, ncclMin);
IMPL_COLL2(ncclReduceScatter, max, FuncMax, ncclCollReduceScatter, ncclMax);
src/collectives/device/reduce_scatter.h
+6 -1
View File
@@ -1,5 +1,6 @@
/*************************************************************************
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
@@ -9,9 +10,10 @@
#include "collectives.h"
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceScatterRingKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int nthreads = blockDim.x - 1;
const int nthreads = blockDim.x;
const int bid = args->bid;
struct ncclDevComm* comm = args->comm;
struct ncclChannel* channel = comm->channels+blockIdx.x;
@@ -62,9 +64,11 @@ __device__ void ncclReduceScatterRingKernel(struct CollectiveArgs* args) {
}
template<int UNROLL, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceScatterTreeKernel(struct CollectiveArgs* args) { }
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceScatterRingLLKernel(struct CollectiveArgs* args) {
const int tid = threadIdx.x;
const int bid = args->bid;
@@ -120,4 +124,5 @@ __device__ void ncclReduceScatterRingLLKernel(struct CollectiveArgs* args) {
}
template<int UNUSED, class FUNC, typename T>
__attribute__((noinline))
__device__ void ncclReduceScatterTreeLLKernel(struct CollectiveArgs* args) { }