为什么CUDA约简算法跳过数组的整个部分？

Question

我正在尝试使用在官方数据自动化系统减少PDF 这里中提到的还原内核之一。然而，我并不理解它是如何工作的，除非我遗漏了一些似乎不太有意义的东西。

这是我的内核：

__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个ints的数组，并且我启动了内核，每个内核有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();

块0的线程0当然会访问

0* 0*2 +0=0

但是块1的线程0将访问

1* 1024*2 +0= 2048

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

我是不是漏掉了一些很明显的东西？

Answer 1

你有

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

因此，对于线程0块0来说，g_idata[i+blockDim.x]是g_idata[0 + 1024]，因此不会跳过，对吗？