CUDA 入门笔记-2 -并行规约

与CUDA
C权威编程指南的代码不同，在相邻配对2和交错配对的时候，grid.x可以除以二，因为第一周期就有一半的资源没有用上，确实优化了效率。

简介

问题：对有N个元素的整数数组求和。
当数据量很大的时候，我们可以分组求和，即
１、讲输入向量划分到更小的数据块中
２、所有的数据块并行求部分和
３、对所有的部分和进行求和得到结果
接下来介绍各种并行规约的方法。

相邻配对－１

每一个block负责计算出一个部分和。
首先有两个全局数组，ｇ_idata和ｇ_odata，分别存放需要求和的元素和每个block计算得出的部分和。
规约是就地完成的，也就是说每一步，全局内存ｇ_idata中的值都会被部分和所替代。
计算结构如下图：

_syncthreads保证，进入下一次迭代之前，每个线程的部分和都已经被保存在全局内存当中。
核函数逻辑为，图中相邻两个元素的间隔为跨度，第一行跨度为１，第二行跨度为２…，每一次循环规约结束以后，跨度变成原来的两倍。
第一次循环后，偶数元素被（他自己和右侧奇数元素）的部分和替代。
第二次循环后，每四个元素产生的部分和在相应的部分。
…

核函数代码如下：

// Neighbored Pair Implementation with divergence
__global__ void reduceNeighbored (int *g_idata, int *g_odata, unsigned int n)
{
   
    // set thread ID
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x;

    // boundary check
    if (idx >= n) return;

    // in-place reduction in global memory
    for (int stride = 1; stride < blockDim.x; stride *= 2)
    {
   
        if ((tid % (2 * stride)) == 0)
        {
   
            idata[tid] += idata[tid + stride];
        }

        // synchronize within threadblock
        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

相邻配对－２

相邻配对１中判断是不是需要规约的条件为

if ((tid % (2 * stride)) == 0)
        {
   
            idata[tid]