为什么 CUDA 缩减算法会跳过数组的整个部分？答案

【问题标题】：Why does a CUDA Reduction algorithm skip over whole sections of an array?为什么 CUDA 缩减算法会跳过数组的整个部分？
【发布时间】：2016-03-03 03:48:52
【问题描述】：

我正在尝试使用官方 CUDA Reduction PDF here 中谈到的一个缩减内核。但是，我不明白它是如何工作的，除非我遗漏了一些似乎没有多大意义的东西。

这是我的内核：

__global__ void energyKernel(int nbodies, int *addReduc, int *subReduc, int *inData, int *inData2){
    extern __shared__ int e[];

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

    if (tID < nbodies && (i + blockDim.x) < nbodies){
        e[tID] = inData[i] + inData[i + blockDim.x];
    }
    else{
        e[tID] = inData[i];
    }

    __syncthreads();

    for (unsigned int stride = blockDim.x / 2; stride > 32; stride >>= 1)
    {
        if (tID < stride)
        {
            e[tID] += e[tID + stride];
        }
        __syncthreads();


    }

    if (tID < 32){ warpReduce(e, tID); }

    if (tID == 0)
    {
        addReduc[blockIdx.x] = e[0];
    } 
}

我试图弄清楚将数据从 inData 数组加载到共享内存的工作原理，因此我进行了一些计算。代码原加载数据如下：

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
e[tid] = inData[i];
__syncthreads();

因此，如果 inData 是一个 3000 个整数的数组，并且我启动了带有 3 个块的内核，每个块有 1024 个线程，那么每个线程将使用其线程“地址”访问 inData 数组，即块 1 的线程 512 将是

1 * 1024 + 512 = 1536

所以它会访问 inData 的第 1536 个元素。

但是使用新代码：

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
__syncthreads();

Block 0 的线程 0 当然会访问

0 * 0*2 + 0 = 0

但是 Block 1 的线程 0 可以访问

1 * 1024*2 + 0 = 2048

因此，从 1024 到 2047 的 inData 值将被完全忽略。

我在这里遗漏了一些非常明显的东西吗？

【问题讨论】：

标签： c++ arrays cuda

【解决方案1】：

你有

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];

因此对于线程 0 块 0 你有 g_idata[i+blockDim.x] 是 g_idata[0 + 1024] ，因此它不会被跳过，对吧？

【讨论】：

哇。当然。那是一个愚蠢的问题。谢谢！
这样的细节经常被忽视:)