The reductions so far compute the sum of an array of integers. To
generalize them to a broader set of operations, we turn to C++
templates. blockReduceCG() is already a template on the
reduced type, and cg::reduce() accepts any binary operator
– cg::plus, cg::less (for minimum), a custom
functor, or a lambda – so a warp or block reduction of any type and
operator needs no new machinery.
The templated reduction kernels in the accompanying source code
(reduction1Templated.cuh,
reduction2Templated.cuh, and so on) take two template
parameters, because the output type may differ from the input type:
T is the type being reduced, and ReductionType
is the type used for the partials and the result. The +=
operator “rakes” through the input, accumulating a partial for each
thread, exactly as in the integer kernels but operating on
ReductionType.
A ReductionType that stores its partials in shared
memory must specialize the SharedMemory template, an idiom
from the CUDA SDK that declares the variable-sized shared memory without
provoking alignment errors:
template<class T>
struct SharedMemory
{
__device__ inline operator T*() { extern __shared__ int __smem[]; return (T*) (void *) __smem; }
__device__ inline operator const T*() const { extern __shared__ int __smem[]; return (T*) (void *) __smem; }
};Listing 12-5 shows a class intended to be used with these templated
reductions. CReduction_Sumi_isq computes both the sum and
the sum of squares of an array of integers in a single pass, by defining
operator+= (and the matching SharedMemory
specialization). Reducing to it is a matter of naming it as the
ReductionType.
struct CReduction_Sumi_isq {public: CReduction_Sumi_isq(); int sum; long long sumsq; CReduction_Sumi_isq& operator +=( int a ); volatile CReduction_Sumi_isq& operator +=( int a ) volatile; CReduction_Sumi_isq& operator +=( const CReduction_Sumi_isq& a ); volatile CReduction_Sumi_isq& operator +=( volatile CReduction_Sumi_isq& a ) volatile; }; inline __device__ __host__CReduction_Sumi_isq::CReduction_Sumi_isq(){ sum = 0; sumsq = 0;} inline __device__ __host__CReduction_Sumi_isq&CReduction_Sumi_isq::operator +=( int a ){ sum += a; sumsq += (long long) a*a; return *this;} inline __device__ __host__volatile CReduction_Sumi_isq&CReduction_Sumi_isq::operator +=( int a ) volatile{ sum += a; sumsq += (long long) a*a; return *this;} inline __device__ __host__CReduction_Sumi_isq&CReduction_Sumi_isq::operator +=( const CReduction_Sumi_isq& a ){ sum += a.sum; sumsq += a.sumsq; return *this;} inline __device__ __host__volatile CReduction_Sumi_isq&CReduction_Sumi_isq::operator +=( volatile CReduction_Sumi_isq& a ) volatile{ sum += a.sum; sumsq += a.sumsq; return *this;} inline intoperator!=( const CReduction_Sumi_isq& a, const CReduction_Sumi_isq& b ){ return a.sum != b.sum && a.sumsq != b.sumsq;} //// from Reduction SDK sample:// specialize to avoid unaligned memory // access compile errors//template<>struct SharedMemory<CReduction_Sumi_isq>{ __device__ inline operator CReduction_Sumi_isq*() { extern __shared__ CReduction_Sumi_isq __smem_CReduction_Sumi_isq[]; return (CReduction_Sumi_isq*)__smem_CReduction_Sumi_isq; } __device__ inline operator const CReduction_Sumi_isq*() const { extern __shared__ CReduction_Sumi_isq __smem_CReduction_Sumi_isq[]; return (CReduction_Sumi_isq*)__smem_CReduction_Sumi_isq; }};CReduction_Sumi_isq class. (source on GitHub)The templates in this section fix the reduction’s type and operator
when the program is compiled, and so do the library reductions:
thrust::reduce and thrust::transform_reduce
take their operators as template parameters. A functor can still carry a
runtime value – a threshold, a coefficient – but the operator’s
code and the accumulator type are settled at compile time. When they are
not known until runtime, NVRTC (Section 3.4) closes the gap by compiling
CUDA C++ source into a loadable module while the program runs.
The accompanying reductionNVRTC.cu sample expresses a
reduction as a runtime transform_reduce over a user-supplied
monoid: an accumulator type Acc, a load that
maps an input element to an Acc, a combine
that merges two of them, and an identity. These are exactly
the four pieces CReduction_Sumi_isq supplies at compile
time; here they are strings, substituted into a fixed kernel template
and handed to NVRTC.
struct Monoid {
const char *In; // input element type
const char *AccDecl; // accumulator type
const char *CtxDecl; // optional captured context
const char *loadBody; // In[,Ctx] -> Acc
const char *combineBody; // (Acc,Acc) -> Acc
const char *idBody; // -> Acc
};
// sum and sum of squares, generated at runtime
Monoid sumsq = {
"int", "struct Acc { long long sum, sumsq; };", 0,
"Acc r; r.sum = x; r.sumsq = (long long) x * x; return r;",
"Acc r; r.sum = a.sum + b.sum; r.sumsq = a.sumsq + b.sumsq; return r;",
"Acc r; r.sum = 0; r.sumsq = 0; return r;"
};The generated kernel is the election-based single pass of Section
12.3.2, which fits an arbitrary Acc: the only atomic is the
integer election counter, so the accumulator itself flows only through
combine and never needs a hardware atomic of its own.
Because the pieces are text, one program can generate reductions a
compile-time library could not – the sample builds the
sum-and-sum-of-squares reducer above and, in the same run with no
recompilation, a “count the elements greater than a threshold”
reducer.
The threshold is a captured value, and supplying it changes the
generated kernel’s interface: load gains a
const Ctx& parameter and the kernel a
const Ctx*, so the host assembles a matching launch
argument list. Choosing the election counter as a kernel argument versus
an NVRTC module __device__ global – reached from the host
with cuModuleGetGlobal() – changes the signature the same
way. In each case, the source NVRTC compiles determines the interface,
and the host adapts to it: the flexibility a fixed, precompiled kernel
cannot offer.