CUDA SDK 中的 Reduce3 示例

Question

我正在阅读 CUDA SDK 中的缩减优化，但在从 reduce2 到 reduce3 发生的事情之后我遇到了问题：

/*
    This version uses sequential addressing -- no divergence or bank conflicts.
*/
template <class T>
__global__ void
reduce2(T *g_idata, T *g_odata, unsigned int n)
{
    T *sdata = SharedMemory<T>();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;

    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=blockDim.x/2; s>0; s>>=1)
    {
        if (tid < s)
        {
            sdata[tid] += sdata[tid + s];
        }

        __syncthreads();
    }

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

/*
This version uses n/2 threads --
it performs the first level of reduction when reading from global memory.
*/
template <class T>
__global__ void
reduce3(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
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;

    T mySum = (i < n) ? g_idata[i] : 0;

    if (i + blockDim.x < n)
        mySum += g_idata[i+blockDim.x];

    sdata[tid] = mySum;
    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=blockDim.x/2; s>0; s>>=1)
    {
        if (tid < s)
        {
            sdata[tid] = mySum = mySum + sdata[tid + s];
        }

        __syncthreads();
    }

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

我无法想象 reduce3 的第一级减少尝试做什么，或者为什么线程数减少了一半。谁能指点一下？

Answer 1

两者之间的唯一区别是reduce3在共享内存减少之前执行求和。因此，reduce2 仅从全局内存中加载单个值并将其存储在共享内存中：

// load shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

sdata[tid] = (i < n) ? g_idata[i] : 0;

reduce3 加载 两个 值，将它们相加，然后将结果存储在共享内存中：

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;

T mySum = (i < n) ? g_idata[i] : 0;

if (i + blockDim.x < n)
    mySum += g_idata[i+blockDim.x];

sdata[tid] = mySum;
__syncthreads();

因为标准“2 的幂”减少的第一级是在共享内存减少之前由每个线程完成的，所以所需的线程总数是 reduce2 的一半。您还应该注意，reduce2 中使用的一半线程实际上被浪费了——它们只将数据加载到共享内存中，根本不参与缩减。因此，reduce3 将它们移除并使用更少的线程来执行相同的操作。

如果您继续查看代码的版本，您会看到这个想法扩展到其逻辑结论，每个线程在将结果存储到共享内存并执行归约之前加载并求和许多值。在内存带宽受限的操作中可以提高效率，例如通过使用更少的线程来减少，允许将每个线程的大部分设置开销分摊到更多的输入值上，并减少内存控制器资源的争用。