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

Merge branch 'master' into xgmi_bench

Browse Source
This commit is contained in:
rpathani
2019-08-16 10:56:56 +05:30
committed by GitHub
parent 65e2f5d87b 50c2202fe9
commit deea20d49c
128 changed files with 7416 additions and 6097 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tools/rccl-prim-test/copy_kernel.h
+182 -127
View File
@@ -61,125 +61,6 @@ __device__ inline void ReduceCopy(
}
}
typedef ulong2 Pack128;
template<class FUNC, typename T>
struct MULTI128 {
__device__ void operator()(Pack128& x, Pack128& y) {
x.x = MULTI<FUNC, T>()(x.x, y.x);
x.y = MULTI<FUNC, T>()(x.y, y.y);
}
};
inline __device__ void Fetch128(Pack128& v, Pack128* p) {
v.x = p->x;
v.y = p->y;
}
inline __device__ void Store128(Pack128* p, Pack128& v) {
p->x = v.x;
p->y = v.y;
}
#define WARP_SIZE 32
template<class FUNC, typename T, bool TWO_INPUTS, bool TWO_OUTPUTS, int UNROLL>
__attribute__((noinline))
__device__ inline void ReduceCopy128b( const int w, const int nw, const int t,
Pack128 * src0, Pack128 * src1, Pack128 * dest0, Pack128 * dest1,
const int N) {
Pack128 t0[UNROLL];
Pack128 t1[UNROLL];
const Pack128* src0_end = src0 + N;
const int inc = nw * UNROLL * WARP_SIZE;
const int offset = w * UNROLL * WARP_SIZE + t;
src0 += offset; if (TWO_INPUTS) src1 += offset;
dest0 += offset; if (TWO_OUTPUTS) dest1 += offset;
while (src0 < src0_end) {
#pragma unroll
for (int u = 0; u < UNROLL; ++u) {
Fetch128(t0[u], src0+u*WARP_SIZE);
if (TWO_INPUTS) Fetch128(t1[u], src1+u*WARP_SIZE);
}
#pragma unroll
for (int u = 0; u < UNROLL; ++u) {
if (TWO_INPUTS) MULTI128<FUNC, T>()(t0[u], t1[u]);
Store128(dest0+u*WARP_SIZE, t0[u]);
if (TWO_OUTPUTS) Store128(dest1+u*WARP_SIZE, t0[u]);
}
src0 += inc; if (TWO_INPUTS) src1 += inc;
dest0 += inc; if (TWO_OUTPUTS) dest1 += inc;
}
}
template<int UNROLL, class FUNC, typename T, bool HAS_DEST1, bool HAS_SRC1>
__attribute__((noinline))
__device__ inline void ReduceOrCopy(const int tid, const int nthreads,
volatile T * __restrict__ dest0, volatile T * __restrict__ dest1,
const volatile T * __restrict__ src0, const volatile T * __restrict__ src1,
int N) {
int Nrem = N;
if (Nrem <= 0) return;
int Npreamble = (Nrem<alignof(Pack128)) ? Nrem : AlignUp(dest0, alignof(Pack128)) - dest0;
// stage 0: check if we'll be able to use the fast, 128-bit aligned path.
// If not, we'll just use the slow preamble path for the whole operation
bool alignable = (((AlignUp(src0, alignof(Pack128)) == src0 + Npreamble)) &&
(!HAS_DEST1 || (AlignUp(dest1, alignof(Pack128)) == dest1 + Npreamble)) &&
(!HAS_SRC1 || (AlignUp(src1, alignof(Pack128)) == src1 + Npreamble)));
if (!alignable) {
Npreamble = Nrem;
}
// stage 1: preamble: handle any elements up to the point of everything coming
// into alignment
ReduceCopy<FUNC, T, HAS_SRC1, HAS_DEST1>(tid, nthreads, src0, src1, dest0, dest1, Npreamble);
Nrem -= Npreamble;
if (Nrem == 0) return;
dest0 += Npreamble; if (HAS_DEST1) { dest1 += Npreamble; }
src0 += Npreamble; if (HAS_SRC1) { src1 += Npreamble; }
// stage 2: fast path: use 128b loads/stores to do the bulk of the work,
// assuming the pointers we have are all 128-bit alignable.
int w = tid / WARP_SIZE; // Warp number
int nw = nthreads / WARP_SIZE; // Number of warps
int t = tid % WARP_SIZE; // Thread (inside the warp)
const int PackFactor = sizeof(Pack128) / sizeof(T);
// stage 2a: main loop
int Nalign2a = (Nrem / (PackFactor * UNROLL * nthreads))
* (UNROLL * nthreads); // round down
ReduceCopy128b<FUNC, T, HAS_SRC1, HAS_DEST1, UNROLL>(w, nw, t, (Pack128*)src0, (Pack128*)src1, (Pack128*)dest0, (Pack128*)dest1, Nalign2a);
int Ndone2a = Nalign2a * PackFactor;
Nrem -= Ndone2a;
if (Nrem == 0) return;
dest0 += Ndone2a; if (HAS_DEST1) { dest1 += Ndone2a; }
src0 += Ndone2a; if (HAS_SRC1) { src1 += Ndone2a; }
// stage 2b: slightly less optimized for section when we don't have full
// UNROLLs
int Nalign2b = Nrem / PackFactor;
ReduceCopy128b<FUNC, T, HAS_SRC1, HAS_DEST1, 1>(w, nw, t, (Pack128*)src0, (Pack128*)src1, (Pack128*)dest0, (Pack128*)dest1, Nalign2b);
int Ndone2b = Nalign2b * PackFactor;
Nrem -= Ndone2b;
if (Nrem == 0) return;
dest0 += Ndone2b; if (HAS_DEST1) { dest1 += Ndone2b; }
src0 += Ndone2b; if (HAS_SRC1) { src1 += Ndone2b; }
// stage 2c: tail
ReduceCopy<FUNC, T, HAS_SRC1, HAS_DEST1>(tid, nthreads, src0, src1, dest0, dest1, Nrem);
}
template<typename T>
struct FuncPassA {
__device__ T operator()(const T x, const T y) const {
@@ -217,6 +98,160 @@ struct MULTI<FUNC, float> {
}
};
typedef ulong2 Pack128;
template<class FUNC, typename T>
struct MULTI128 {
__device__ void operator()(Pack128& x, Pack128& y) {
x.x = MULTI<FUNC, T>()(x.x, y.x);
x.y = MULTI<FUNC, T>()(x.y, y.y);
}
};
inline __device__ void Fetch128(Pack128& v, const Pack128* p) {
v.x = p->x;
v.y = p->y;
}
inline __device__ void Store128(Pack128* p, Pack128& v) {
p->x = v.x;
p->y = v.y;
}
template<class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
__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) {
T val = vFetch(srcs[0]+idx);
#pragma unroll
for (int i=1; i<MINSRCS; i++) val = FUNC()(val, vFetch(srcs[i]+idx));
#pragma unroll 1
for (int i=MINSRCS; i<MAXSRCS && i<nsrcs; i++) val = FUNC()(val, vFetch(srcs[i]+idx));
#pragma unroll
for (int i=0; i<MINDSTS; i++) vStore(dsts[i]+idx, val);
#pragma unroll 1
for (int i=MINDSTS; i<MAXDSTS && i<ndsts; i++) vStore(dsts[i]+idx, val);
}
}
#define WARP_SIZE 64
template<class FUNC, typename T, int UNROLL, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
__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;
int offset = w * UNROLL * WARP_SIZE + t;
const Pack128* srcs[MAXSRCS];
for (int i=0; i<MAXSRCS; i++) srcs[i] = ((const Pack128*)(s[i]+elemOffset))+offset;
Pack128* dsts[MAXDSTS];
for (int i=0; i<MAXDSTS; i++) dsts[i] = ((Pack128*)(d[i]+elemOffset))+offset;
while (offset < Npack) {
Pack128 vals[UNROLL];
// Load and reduce
for (int u = 0; u < UNROLL; ++u) Fetch128(vals[u], srcs[0]+u*WARP_SIZE);
for (int i=1; i<MINSRCS; i++) {
Pack128 vals2[UNROLL];
for (int u = 0; u < UNROLL; ++u) Fetch128(vals2[u], srcs[i]+u*WARP_SIZE);
for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>()(vals[u], vals2[u]);
}
#pragma unroll 1
for (int i=MINSRCS; i<MAXSRCS && i<nsrcs; i++) {
Pack128 vals2[UNROLL];
for (int u = 0; u < UNROLL; ++u) Fetch128(vals2[u], srcs[i]+u*WARP_SIZE);
for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>()(vals[u], vals2[u]);
}
// Store
for (int i = 0; i < MINDSTS; i++) {
for (int u = 0; u < UNROLL; ++u) Store128(dsts[i]+u*WARP_SIZE, vals[u]);
}
#pragma unroll 1
for (int i=MINDSTS; i<MAXDSTS && i<ndsts; i++) {
for (int u = 0; u < UNROLL; ++u) Store128(dsts[i]+u*WARP_SIZE, vals[u]);
}
for (int i=0; i<MAXSRCS; i++) srcs[i] += inc;
for (int i=0; i<MAXDSTS; i++) dsts[i] += inc;
offset += inc;
}
}
template <typename T>
__device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(Pack128); }
// Try to limit consecutive load/stores to 8.
// Use UNROLL 8 when we have a single source and a single destination, 4 otherwise
#define AUTOUNROLL (UNROLL*(4/(MINDSTS+MINSRCS)))
template<int UNROLL, class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
__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;
if (Nrem <= 0) return;
int alignDiff = 0;
int align = ptrAlign128(srcs[0]);
#pragma unroll
for (int i=1; i<MINSRCS; i++) alignDiff |= (align ^ ptrAlign128(srcs[i]));
for (int i=MINSRCS; i<MAXSRCS && i<nsrcs; i++) alignDiff |= (align ^ ptrAlign128(srcs[i]));
#pragma unroll
for (int i=0; i<MINDSTS; i++) alignDiff |= (align ^ ptrAlign128(dsts[i]));
for (int i=MINDSTS; i<MAXDSTS && i<ndsts; i++) alignDiff |= (align ^ ptrAlign128(dsts[i]));
int Npreamble = alignDiff ? Nrem :
N < alignof(Pack128) ? N :
(alignof(Pack128) - align) % alignof(Pack128);
// stage 1: preamble: handle any elements up to the point of everything coming
// into alignment
if (Npreamble) {
ReduceCopyMulti<FUNC, T, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(tid, nthreads, nsrcs, srcs, ndsts, dsts, 0, Npreamble);
Nrem -= Npreamble;
if (Nrem == 0) return;
}
int offset = Npreamble;
// stage 2: fast path: use 128b loads/stores to do the bulk of the work,
// assuming the pointers we have are all 128-bit alignable.
int w = tid / WARP_SIZE; // Warp number
int nw = nthreads / WARP_SIZE; // Number of warps
int t = tid % WARP_SIZE; // Thread (inside the warp)
const int packFactor = sizeof(Pack128) / sizeof(T);
// stage 2a: main loop
int Npack2a = (Nrem / (packFactor * AUTOUNROLL * WARP_SIZE))
* (AUTOUNROLL * WARP_SIZE); // round down
int Nelem2a = Npack2a * packFactor;
ReduceCopy128bMulti<FUNC, T, AUTOUNROLL, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Npack2a);
Nrem -= Nelem2a;
if (Nrem == 0) return;
offset += Nelem2a;
// stage 2b: slightly less optimized for section when we don't have full
// unrolling
int Npack2b = Nrem / packFactor;
int Nelem2b = Npack2b * packFactor;
ReduceCopy128bMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Npack2b);
Nrem -= Nelem2b;
if (Nrem == 0) return;
offset += Nelem2b;
// stage 2c: tail
ReduceCopyMulti<FUNC, T, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(tid, nthreads, nsrcs, srcs, ndsts, dsts, offset, Nrem);
}
// Assumptions:
// - there is exactly 1 block
// - THREADS is the number of producer threads
@@ -224,24 +259,38 @@ struct MULTI<FUNC, float> {
template<int UNROLL, int THREADS, typename T>
__device__ void Copy(volatile T * __restrict__ const dest,
const volatile T * __restrict__ const src, const int N) {
ReduceOrCopy<UNROLL, FuncPassA<T>, T, false, false>(threadIdx.x, THREADS,
dest, nullptr, src, nullptr, N);
const T* srcs[2];
T* dsts[2];
srcs[0] = (const T*)src;
dsts[0] = (T*)dest;
ReduceOrCopyMulti<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
1, srcs, 1, dsts, N);
}
template<int UNROLL, int THREADS, typename T>
__device__ void DoubleCopy(volatile T * __restrict__ const dest0,
volatile T * __restrict__ const dest1,
const volatile T * __restrict__ const src, const int N) {
ReduceOrCopy<UNROLL, FuncPassA<T>, T, true, false>(threadIdx.x, THREADS,
dest0, dest1, src, nullptr, N);
const T* srcs[2];
T* dsts[2];
srcs[0] = (const T*)src;
dsts[0] = (T*)dest0;
dsts[1] = (T*)dest1;
ReduceOrCopyMulti<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
1, srcs, 2, dsts, N);
}
template<int UNROLL, int THREADS, typename T>
__device__ void Reduce(volatile T * __restrict__ const dest,
const volatile T * __restrict__ const src0,
const volatile T * __restrict__ const src1, const int N) {
ReduceOrCopy<UNROLL, FuncSum<T>, T, false, true>(threadIdx.x, THREADS,
dest, nullptr, src0, src1, N);
const T* srcs[2];
T* dsts[2];
srcs[0] = (const T*)src0;
srcs[1] = (const T*)src1;
dsts[0] = (T*)dest;
ReduceOrCopyMulti<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
2, srcs, 1, dsts, N);
}
template<int UNROLL, int THREADS, typename T>
@@ -249,7 +298,13 @@ __device__ void ReduceCopy(volatile T * __restrict__ const dest0,
volatile T * __restrict__ const dest1,
const volatile T * __restrict__ const src0,
const volatile T * __restrict__ const src1, const int N) {
ReduceOrCopy<UNROLL, FuncSum<T>, T, true, true>(threadIdx.x, THREADS,
dest0, dest1, src0, src1, N);
const T* srcs[2];
T* dsts[2];
srcs[0] = (const T*)src0;
srcs[1] = (const T*)src1;
dsts[0] = (T*)dest0;
dsts[1] = (T*)dest1;
ReduceOrCopyMulti<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
2, srcs, 2, dsts, N);
}
#endif // COPY_KERNEL_H_
tools/rccl-prim-test/rccl_prim_test.cpp
+4 -2
View File
@@ -427,10 +427,12 @@ int main(int argc,char* argv[])
sizeof(struct profiling_data_t), hipMemcpyDeviceToHost,
stream[i]));
HIPCHECK(hipStreamSynchronize(stream[i]));
int next_gpu = findNextGpu(ring_0, i, nGpu);
uint32_t linktype;
uint32_t hopcount;
HIPCHECK(hipExtGetLinkTypeAndHopCount(i, next_gpu , &linktype, &hopcount));
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, i));
if(prop.gcnArch == 906 ) {
@@ -441,11 +443,11 @@ int main(int argc,char* argv[])
double t0 = (double)profiling_data[i]->write_cycles/((double)RTC_CLOCK_FREQ_MI100)/(double)workgroups;
fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n",
i, next_gpu,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9));
} else {
} else {
double t0 = (double)profiling_data[i]->write_cycles/((double)RTC_CLOCK_FREQ_DEFAULT)/(double)workgroups;
fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n",
i, next_gpu,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9));
}
}
}
std::cout<<"***Application Level Transfer Profiling Data***"<<std::endl;
double speed = (double)(profiling_data[0]->bytes_transferred) / (deltaSec*1.0E9);