c++ - 为什么 CUDA Reduction 算法会跳过数组的整个部分？

Question

我正在尝试使用官方 CUDA 缩减 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 个 block 的内核，每个 block 具有 1024 个线程，那么每个线程将使用其线程“地址”访问 inData 数组，即 block 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 的 Thread 0 当然会访问

0 * 0*2 + 0 = 0

但是 block 1 的线程 0 会访问

1 * 1024*2 + 0 = 2048

因此完全忽略 1024 - 2047 的 inData 值。

我是否遗漏了一些非常明显的东西？

Answer 1

你有

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

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