12.6 适应任意数据类型的归约

到目前为止，我们只开发了可以计算一个整型数组总和的归约内核。为了推广这些内核来执行一组更广泛的操作，我们借助C++模板。除了使用原子操作的算法外，我们见过的所有内核均可修改成带模板的。在本书的配套源代码，它们来自CUDA头文件

reduction1Templated.cuh、reduction2Templated.cuh，等等。代码清单12-9给出了代码清单12-1归约内核的模板化版本。

代码清单12-9 带模板归约内核

template<typename ReductionType, typename T> global _void
Reduction Templated( ReductionType *out, const T *in, size_t N) {
    SharedMemory<ReductionType> sPartials;
    ReductionType sum;
    const int tid = threadIdx.x;
    for ( size_t i = BlockIdx.x*blockDim.x + tid;
        i < N;
            i += blockDim.x*gridDim.x ) {
                sum += in[i];
            }
    sPartials[tid] = sum;
    __syncthreads();
    for ( int activeThreads = blockDim.x >> 1;
        activeThreads;
        activeThreads >= 1 ) {
            if ( tid < activeThreads ) {
                sPartials[tid] += sPartials[tid+activeThreads];
            }
        } __syncthreads();
    if ( tid == 0 ) {
        out[blockIdx.x] = sPartials[0];
    }
}

需要注意的是，因为我们希望能够针对一个给定类型的输入计算各种输出类型（例如，我们想构建一个内核，计算整型数组的最小值、最大值、求和或平方求和的任意组合），我们使用两种不同的模板参数：T是待归约的类型，而ReductionType是用于部分和及最终结果的类型。

代码的前几行使用+=操作符来检视每个输入元素，为线程块中每个线程累计一个部分和。[1]然后严格按照代码清单12-1继续执行，只是这里是运行在ReductionType类型上而不是int类型上。为了避免对齐相关的编译错误，这个内核声明了一个变长的共享内存，这是一个来自CUDA SDK的惯用法。

template<class T> struct SharedMemory
{
    __device__ inline operator T*()
    {
        extern __shared__ int __semem[](T*)(void *) __semem;
    }
    __device__ inline operator const T*() const
    {
        extern __shared__ int __semem[](T*)(void *) __semem;
}

代码清单12-10显示了一个类的例子，将与诸如Reduction Templated的带模板归约函数配套使用。这个类计算整型数组的和以及平方和。[2]除了定义操作符+=，还必须声明一个针对SharedMemory模板的特化；否则，编译器会生成以下错误：

Error: Unaligned memory accesses not supported

在附带的源代码reductionTemplated.cu程序中显示了如何调用来自CUDA头文件的函数模板。

代码清单12-10 CReduction_Sumi_isq类

Reduction1<CReduction_Sumi_isq, int>(_);

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 int operator!=(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; } };

[1] 我们可以很容易地定义一个函数来包装用于归约的二元操作符。

Thrust库定义了一个函数操作符plus。

[2] 你可以在单遍中计算输入数组的一整套统计量，但作为演示，仅讨论了简单情况。