用CUDA进行求和缩减：N是什么？

Question

据NVIDIA公司称，这个是最快的加和降维核心。

template <unsigned int blockSize>
__device__ void warpReduce(volatile int *sdata, unsigned int tid) {
if (blockSize >=  64) sdata[tid] += sdata[tid + 32];
if (blockSize >=  32) sdata[tid] += sdata[tid + 16];
if (blockSize >=  16) sdata[tid] += sdata[tid +  8];
if (blockSize >=    8) sdata[tid] += sdata[tid +  4];
if (blockSize >=    4) sdata[tid] += sdata[tid +  2];
if (blockSize >=    2) sdata[tid] += sdata[tid +  1];
}
template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n) {
extern __shared__ int sdata[];
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockSize*2) + tid;
unsigned int gridSize = blockSize*2*gridDim.x;
sdata[tid] = 0;
while (i < n) { sdata[tid] += g_idata[i] + g_idata[i+blockSize];  i += gridSize;  }
__syncthreads();
if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
if (blockSize >= 128) { if (tid <   64) { sdata[tid] += sdata[tid +   64]; } __syncthreads(); }
if (tid < 32) warpReduce(sdata, tid);
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

然而，我不理解“n”参数的含义。有什么线索吗？我不认为它是要缩小的数组大小，因为在while循环中可能会发生缓冲区溢出。

- Michael Eilers Smith

---

n是g_idata数组中元素的数量。此外，那个特定的内核不太可能是“最快”的约简内核；这个文档现在已经相当老了。 - Jared Hoberock

请注意，使用NVidia Kepler架构，您引用的代码肯定不是最快的约简方式。可以使用_shfl_xor指令完成intra-warp工作。请参阅此演示文稿：http://on-demand.gputechconf.com/gtc/2013/presentations/S3174-Kepler-Shuffle-Tips-Tricks.pdf。 - einpoklum

Answer 1

我相信您在幻灯片中发现了一个错别字（应该是类似于while(i + blockDim.x < n)的内容）。

如果您查看CUDA SDK示例"reduction"中的源代码，最新的reduce6主体如下：

template <class T, unsigned int blockSize, bool nIsPow2>
__global__ void
reduce6(T *g_idata, T *g_odata, unsigned int n)
{
    T *sdata = SharedMemory<T>();

    // perform first level of reduction,
    // reading from global memory, writing to shared memory
    ...

    T mySum = 0;

    // we reduce multiple elements per thread.  The number is determined by the 
    // number of active thread blocks (via gridDim).  More blocks will result
    // in a larger gridSize and therefore fewer elements per thread
    while (i < n)
    {         
        mySum += g_idata[i];
        // ensure we don't read out of bounds -- this is optimized away for powerOf2 sized arrays
        if (nIsPow2 || i + blockSize < n) 
            mySum += g_idata[i+blockSize];  
        i += gridSize;
    } 

请注意while中的显式检查，以防止对g_idata进行越界访问。您最初的怀疑是正确的; n只是g_idata数组的大小。

- Jared Hoberock