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

Merge remote-tracking branch 'nccl/master' into 2.10.3

Browse Source
This commit is contained in:
Wenkai Du
2021-07-28 13:27:06 -07:00
parent 229ca88ee6 7e51592129
commit bf2339f93e
53 changed files with 3806 additions and 3023 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/collectives/device/common_kernel.h
+314 -77
View File
@@ -1,5 +1,5 @@
/*************************************************************************
* Copyright (c) 2015-2020, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2015-2021, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
@@ -17,24 +17,26 @@
// Define min for ssize_t
static __device__ int min(int a, ssize_t b) { return (a < b) ? a : b; }
template <typename T>
inline __device__ void loadPtr(void** ptr, T* &v) {
v = LOAD(ptr);
}
typedef uint64_t PackType;
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
template<class FUNC, typename T>
struct MULTI {
__device__ PackType operator()(const PackType x, const PackType y) const
{
return FUNC()(x, y);
}
};
#else
template<typename Fn>
struct FuncTraits /*{
__device__ static Fn make();
__device__ static T preOp(Fn, T);
__device__ static T postOp(Fn, T);
}*/;
// unpack x and y to elements of type T and apply FUNC to each element
template<class FUNC, typename T>
struct MULTI {
__device__ PackType operator()(const PackType x, const PackType y) const;
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const;
__device__ PackType preOp(FUNC fn, PackType x) const;
__device__ PackType postOp(FUNC fn, PackType x) const;
};
template<class FUNC>
@@ -48,17 +50,39 @@ struct MULTI<FUNC, int8_t> {
};
};
__device__ PackType operator()(const PackType x, const PackType y) const {
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
converter cx, cy, cr;
cx.storage = x;
cy.storage = y;
// for char, we do these as vector ops
cr.a = FUNC()(cx.a, cy.a);
cr.b = FUNC()(cx.b, cy.b);
cr.a = fn(cx.a, cy.a);
cr.b = fn(cx.b, cy.b);
return cr.storage;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
int8_t elt[8];
} u;
u.pack = x;
#pragma unroll 1
for (int i=0; i < 8; i++)
u.elt[i] = FuncTraits<FUNC>().preOp(fn, u.elt[i]);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
int8_t elt[8];
} u;
u.pack = x;
#pragma unroll 1
for (int i=0; i < 8; i++)
u.elt[i] = FuncTraits<FUNC>().postOp(fn, u.elt[i]);
return u.pack;
}
};
template<class FUNC>
@@ -72,17 +96,39 @@ struct MULTI<FUNC, uint8_t> {
};
};
__device__ PackType operator()(const PackType x, const PackType y) const {
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
converter cx, cy, cr;
cx.storage = x;
cy.storage = y;
// for char, we do these as vector ops
cr.a = FUNC()(cx.a, cy.a);
cr.b = FUNC()(cx.b, cy.b);
cr.a = fn(cx.a, cy.a);
cr.b = fn(cx.b, cy.b);
return cr.storage;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
uint8_t elt[8];
} u;
u.pack = x;
#pragma unroll 1
for (int i=0; i < 8; i++)
u.elt[i] = FuncTraits<FUNC>().preOp(fn, u.elt[i]);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
uint8_t elt[8];
} u;
u.pack = x;
#pragma unroll 1
for (int i=0; i < 8; i++)
u.elt[i] = FuncTraits<FUNC>().postOp(fn, u.elt[i]);
return u.pack;
}
};
template<class FUNC>
@@ -96,16 +142,36 @@ struct MULTI<FUNC, int32_t> {
};
};
__device__ PackType operator()(const PackType x, const PackType y) const {
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
converter cx, cy, cr;
cx.storage = x;
cy.storage = y;
cr.a = FUNC()(cx.a, cy.a);
cr.b = FUNC()(cx.b, cy.b);
cr.a = fn(cx.a, cy.a);
cr.b = fn(cx.b, cy.b);
return cr.storage;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
int32_t elt[2];
} u;
u.pack = x;
u.elt[0] = FuncTraits<FUNC>().preOp(fn, u.elt[0]);
u.elt[1] = FuncTraits<FUNC>().preOp(fn, u.elt[1]);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
int32_t elt[2];
} u;
u.pack = x;
u.elt[0] = FuncTraits<FUNC>().postOp(fn, u.elt[0]);
u.elt[1] = FuncTraits<FUNC>().postOp(fn, u.elt[1]);
return u.pack;
}
};
template<class FUNC>
@@ -119,16 +185,36 @@ struct MULTI<FUNC, uint32_t> {
};
};
__device__ PackType operator()(const PackType x, const PackType y) const {
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
converter cx, cy, cr;
cx.storage = x;
cy.storage = y;
cr.a = FUNC()(cx.a, cy.a);
cr.b = FUNC()(cx.b, cy.b);
cr.a = fn(cx.a, cy.a);
cr.b = fn(cx.b, cy.b);
return cr.storage;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
uint32_t elt[2];
} u;
u.pack = x;
u.elt[0] = FuncTraits<FUNC>().preOp(fn, u.elt[0]);
u.elt[1] = FuncTraits<FUNC>().preOp(fn, u.elt[1]);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
uint32_t elt[2];
} u;
u.pack = x;
u.elt[0] = FuncTraits<FUNC>().postOp(fn, u.elt[0]);
u.elt[1] = FuncTraits<FUNC>().postOp(fn, u.elt[1]);
return u.pack;
}
};
template<class FUNC>
@@ -136,22 +222,75 @@ struct MULTI<FUNC, half> {
static_assert(sizeof(PackType) == 4 * sizeof(half),
"PackType must be four times the size of half.");
struct PackHalf2 {
half2 a, b;
union Converter {
PackType pack;
half2 h2[2];
};
__device__ PackType operator()(const PackType x, const PackType y) const {
struct PackHalf2 cx, cy, cr;
cx = *(reinterpret_cast<const struct PackHalf2*>(&x));
cy = *(reinterpret_cast<const struct PackHalf2*>(&y));
cr.a = FUNC()(cx.a, cy.a);
cr.b = FUNC()(cx.b, cy.b);
return *(reinterpret_cast<PackType*>(&cr));
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
Converter cx, cy, cr;
cx.pack = x;
cy.pack = y;
cr.h2[0] = fn(cx.h2[0], cy.h2[0]);
cr.h2[1] = fn(cx.h2[1], cy.h2[1]);
return cr.pack;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
Converter c;
c.pack = x;
c.h2[0] = FuncTraits<FUNC>().preOp(fn, c.h2[0]);
c.h2[1] = FuncTraits<FUNC>().preOp(fn, c.h2[1]);
return c.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
Converter c;
c.pack = x;
c.h2[0] = FuncTraits<FUNC>().postOp(fn, c.h2[0]);
c.h2[1] = FuncTraits<FUNC>().postOp(fn, c.h2[1]);
return c.pack;
}
};
#if defined(RCCL_BFLOAT16)
template<class FUNC>
struct MULTI<FUNC, rccl_bfloat16> {
static_assert(sizeof(PackType) == 4 * sizeof(rccl_bfloat16),
"PackType must be four times the size of rccl_bfloat16.");
union Converter {
PackType pack;
rccl_bfloat16 h2[4];
};
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
Converter cx, cy, cr;
cx.pack = x;
cy.pack = y;
cr.h2[0] = fn(cx.h2[0], cy.h2[0]);
cr.h2[1] = fn(cx.h2[1], cy.h2[1]);
cr.h2[2] = fn(cx.h2[2], cy.h2[2]);
cr.h2[3] = fn(cx.h2[3], cy.h2[3]);
return cr.pack;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
Converter c;
c.pack = x;
c.h2[0] = FuncTraits<FUNC>().preOp(fn, c.h2[0]);
c.h2[1] = FuncTraits<FUNC>().preOp(fn, c.h2[1]);
c.h2[2] = FuncTraits<FUNC>().preOp(fn, c.h2[2]);
c.h2[3] = FuncTraits<FUNC>().preOp(fn, c.h2[3]);
return c.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
Converter c;
c.pack = x;
c.h2[0] = FuncTraits<FUNC>().postOp(fn, c.h2[0]);
c.h2[1] = FuncTraits<FUNC>().postOp(fn, c.h2[1]);
c.h2[2] = FuncTraits<FUNC>().postOp(fn, c.h2[2]);
c.h2[3] = FuncTraits<FUNC>().postOp(fn, c.h2[3]);
return c.pack;
}
};
#endif
template<class FUNC>
struct MULTI<FUNC, float> {
static_assert(sizeof(PackType) == 2 * sizeof(float),
@@ -163,50 +302,122 @@ struct MULTI<FUNC, float> {
};
};
__device__ PackType operator()(const PackType x, const PackType y) const {
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
converter cx, cy, cr;
cx.storage = x;
cy.storage = y;
cr.a = FUNC()(cx.a, cy.a);
cr.b = FUNC()(cx.b, cy.b);
cr.a = fn(cx.a, cy.a);
cr.b = fn(cx.b, cy.b);
return cr.storage;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
float elt[2];
} u;
u.pack = x;
u.elt[0] = FuncTraits<FUNC>().preOp(fn, u.elt[0]);
u.elt[1] = FuncTraits<FUNC>().preOp(fn, u.elt[1]);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
float elt[2];
} u;
u.pack = x;
u.elt[0] = FuncTraits<FUNC>().postOp(fn, u.elt[0]);
u.elt[1] = FuncTraits<FUNC>().postOp(fn, u.elt[1]);
return u.pack;
}
};
template<class FUNC>
struct MULTI<FUNC, double> {
static_assert(sizeof(PackType) == sizeof(double),
"PackType must be the same size as double.");
__device__ PackType operator()(const PackType x, const PackType y) const {
double rv = FUNC()(__longlong_as_double(x), __longlong_as_double(y));
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
double rv = fn(__longlong_as_double(x), __longlong_as_double(y));
return __double_as_longlong(rv);
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
double elt;
} u;
u.pack = x;
u.elt = FuncTraits<FUNC>().preOp(fn, u.elt);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
double elt;
} u;
u.pack = x;
u.elt = FuncTraits<FUNC>().postOp(fn, u.elt);
return u.pack;
}
};
template<class FUNC>
struct MULTI<FUNC, uint64_t> {
static_assert(sizeof(PackType) == sizeof(uint64_t),
"PackType must be the same size as uint64_t.");
__device__ PackType operator()(const PackType x, const PackType y) const {
uint64_t rv = FUNC()(x, y);
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
uint64_t rv = fn(x, y);
return rv;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
uint64_t elt;
} u;
u.pack = x;
u.elt = FuncTraits<FUNC>().preOp(fn, u.elt);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
uint64_t elt;
} u;
u.pack = x;
u.elt = FuncTraits<FUNC>().postOp(fn, u.elt);
return u.pack;
}
};
template<class FUNC>
struct MULTI<FUNC, int64_t> {
static_assert(sizeof(PackType) == sizeof(int64_t),
"PackType must be the same size as int64_t.");
__device__ PackType operator()(const PackType x, const PackType y) const {
int64_t rv = FUNC()((int64_t)x, (int64_t)y);
__device__ PackType operator()(FUNC fn, const PackType x, const PackType y) const {
int64_t rv = fn((int64_t)x, (int64_t)y);
return rv;
}
__device__ PackType preOp(FUNC fn, PackType x) const {
union {
PackType pack;
int64_t elt;
} u;
u.pack = x;
u.elt = FuncTraits<FUNC>().preOp(fn, u.elt);
return u.pack;
}
__device__ PackType postOp(FUNC fn, PackType x) const {
union {
PackType pack;
int64_t elt;
} u;
u.pack = x;
u.elt = FuncTraits<FUNC>().postOp(fn, u.elt);
return u.pack;
}
};
#endif //defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
template<typename T> inline __device__
T vFetch(const volatile T* ptr) {
return *ptr;
@@ -259,9 +470,17 @@ 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);
__device__ void operator()(FUNC fn, Pack128& x, Pack128 const& y) const {
x.x = MULTI<FUNC, T>()(fn, x.x, y.x);
x.y = MULTI<FUNC, T>()(fn, x.y, y.y);
}
__device__ void preOp(FUNC fn, Pack128 &x) const {
x.x = MULTI<FUNC, T>().preOp(fn, x.x);
x.y = MULTI<FUNC, T>().preOp(fn, x.y);
}
__device__ void postOp(FUNC fn, Pack128 &x) const {
x.x = MULTI<FUNC, T>().postOp(fn, x.x);
x.y = MULTI<FUNC, T>().postOp(fn, x.y);
}
};
@@ -284,7 +503,8 @@ inline __device__ void Store128(Pack128* p, Pack128& v) {
template<class FUNC, typename T, int UNROLL, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
__device__ __forceinline__ void ReduceCopyMulti(const int w, const int nw, const int t,
int nsrcs, const T** s, int ndsts, T** d, const int elemOffset, const int Nelem) {
FUNC fn, bool preOpSrc0, bool postOp, int nsrcs, const T** s, int ndsts, T** d, const int elemOffset, const int Nelem
) {
const int inc = nw * UNROLL * WARP_SIZE;
int offset = w * UNROLL * WARP_SIZE + t;
@@ -297,22 +517,30 @@ __device__ __forceinline__ void ReduceCopyMulti(const int w, const int nw, const
T vals[UNROLL];
// Load and reduce
for (int u = 0; u < UNROLL; ++u) vals[u] = vFetch(srcs[0]+u*WARP_SIZE);
if (preOpSrc0) {
for (int u = 0; u < UNROLL; ++u) vals[u] = FuncTraits<FUNC>().preOp(fn, vals[u]);
}
#pragma unroll
for (int i=1; i<MINSRCS; i++) {
T vals2[UNROLL];
for (int u = 0; u < UNROLL; ++u) vals2[u] = vFetch(srcs[i]+u*WARP_SIZE);
for (int u = 0; u < UNROLL; ++u) vals[u] = FUNC()(vals[u], vals2[u]);
for (int u = 0; u < UNROLL; ++u) vals[u] = fn(vals[u], vals2[u]);
}
#pragma unroll
for (int i=MINSRCS; i<MAXSRCS; i++) {
if (i<nsrcs) {
T vals2[UNROLL];
for (int u = 0; u < UNROLL; ++u) vals2[u] = vFetch(srcs[i]+u*WARP_SIZE);
for (int u = 0; u < UNROLL; ++u) vals[u] = FUNC()(vals[u], vals2[u]);
for (int u = 0; u < UNROLL; ++u) vals[u] = fn(vals[u], vals2[u]);
}
}
if (postOp) {
#pragma unroll
for (int u = 0; u < UNROLL; ++u) vals[u] = FuncTraits<FUNC>().postOp(fn, vals[u]);
}
// Store
#pragma unroll
for (int i = 0; i < MINDSTS; i++) {
@@ -331,8 +559,9 @@ __device__ __forceinline__ void ReduceCopyMulti(const int w, const int nw, const
}
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, int ndsts, T** d, const int elemOffset, const int Npack) {
__device__ __forceinline__ void ReduceCopy128bMulti(const int w, const int nw, const int t,
FUNC fn, bool preOpSrc0, bool postOp, int nsrcs, const T** s, int ndsts, T** d, const int elemOffset, const int Npack
) {
const int inc = nw * UNROLL * WARP_SIZE;
int offset = w * UNROLL * WARP_SIZE + t;
@@ -345,20 +574,32 @@ __device__ void ReduceCopy128bMulti(const int w, const int nw, const int t,
Pack128 vals[UNROLL];
// Load and reduce
for (int u = 0; u < UNROLL; ++u) Fetch128(vals[u], srcs[0]+u*WARP_SIZE);
if (preOpSrc0) {
for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>().preOp(fn, vals[u]);
}
#pragma unroll 1
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]);
for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>()(fn, 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]);
for (int i=MINSRCS; i<MAXSRCS; i++) {
if (i<nsrcs) {
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>()(fn, vals[u], vals2[u]);
}
}
if (postOp) {
#pragma unroll 1
for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>().postOp(fn, vals[u]);
}
// Store
#pragma unroll 1
for (int i = 0; i < MINDSTS; i++) {
for (int u = 0; u < UNROLL; ++u) Store128(dsts[i]+u*WARP_SIZE, vals[u]);
}
@@ -375,23 +616,15 @@ __device__ void ReduceCopy128bMulti(const int w, const int nw, const int t,
}
template <typename T>
__device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(int32_t); }
__device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(uint32_t); }
#define PACKELEMS (sizeof(Pack128) / sizeof(T))
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
// Multiply UNROLL by 2 if single source/single destination
#define AUTOUNROLL (UNROLL*((MINSRCS==1 && MINDSTS==1) ? 2 : 1))
#else
// 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)))
#endif
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,
int nsrcs, const T** srcs, int ndsts, T** dsts,
int N) {
__device__ __forceinline__ void ReduceOrCopyMulti(
const int tid, const int nthreads, FUNC fn, bool preOpSrc0, bool postOp, int nsrcs, const T** srcs, int ndsts, T** dsts, int N
) {
int Nrem = N;
if (Nrem <= 0) return;
@@ -417,7 +650,8 @@ __device__ __forceinline__ void ReduceOrCopyMulti(const int tid, const int nthre
int Npack = (Nrem / (PACKELEMS*AUTOUNROLL*WARP_SIZE)) * (AUTOUNROLL*WARP_SIZE); // round down
int Nelem = Npack * PACKELEMS;
ReduceCopy128bMulti<FUNC, T, AUTOUNROLL, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Npack);
ReduceCopy128bMulti<FUNC, T, AUTOUNROLL, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
(w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Npack);
Nrem -= Nelem;
if (Nrem == 0) return;
@@ -427,7 +661,8 @@ __device__ __forceinline__ void ReduceOrCopyMulti(const int tid, const int nthre
Npack = Nrem / PACKELEMS;
Nelem = Npack * PACKELEMS;
ReduceCopy128bMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Npack);
ReduceCopy128bMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
(w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Npack);
Nrem -= Nelem;
if (Nrem == 0) return;
@@ -437,14 +672,16 @@ __device__ __forceinline__ void ReduceOrCopyMulti(const int tid, const int nthre
// unrolled, by-type (mostly for unaligned buffers)
int Nelem = (Nrem / (UNROLL*PACKELEMS/2*WARP_SIZE)) * (UNROLL*PACKELEMS/2*WARP_SIZE); // round down
ReduceCopyMulti<FUNC, T, UNROLL*PACKELEMS/2, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Nelem);
ReduceCopyMulti<FUNC, T, UNROLL*PACKELEMS/2, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
(w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Nelem);
Nrem -= Nelem;
if (Nrem == 0) return;
offset += Nelem;
// no unroll, by type. Should finish what's remaining.
ReduceCopyMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Nrem);
ReduceCopyMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
(w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Nrem);
}
#endif // COMMON_KERNEL_H_