c++ - “官方”CUDA Reduction 函数不能接受某些数字？

Question

目前正在尝试使用 CUDA pdf 中的 Reduction #3 outline here .

这是我的 Reduction 函数的样子

template <typename T>
__device__ void offsetReduction(planet<T> *bodies, T *outdata, int arrayIdent, int nbodies){
    extern __shared__ T sdata[];

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

    if (arrayIdent == 1){
        if (i < nbodies){
            sdata[tID] = bodies[i].vx * bodies[i].mass;
        }

        __syncthreads();
    }

    if (arrayIdent == 2){
        if (i < nbodies){
            sdata[tID] = (bodies[i].vy * bodies[i].mass);
        }
        __syncthreads();
    }

    if (arrayIdent == 3){
        if (i < nbodies){
            sdata[tID] = (bodies[i].vz * bodies[i].mass);
        }
        __syncthreads();
    }

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

    if (tID == 0)
    {
        outdata[blockIdx.x] = sdata[0];
    }

但是，它似乎并没有正常工作，所以我做了一些计算。

我启动了与“int nbodies”相同数量的线程，在我的例子中我选择了 5 个。所以 5 个线程中的每一个都进入并向 sdata[] 添加一个值没问题。然而，一旦它到达加法部分，它就出错了。

在第一次迭代中，线程 0 访问 sdata[3]，线程 1 访问 sdata[4]，其他线程什么都不做。在第二次迭代中，线程 0 访问 sdata 1而其他线程什么都不做。添加然后“完成”并且内核完成。但是 sdata[2] 从未被添加，所以我在 sdata[0] 中得到了一个不正确的值。

我是否漏掉了一些非常明显的东西？ (我已经盯着这个看了一段时间，所以我可能已经看过了。

Answer 1

与任何其他“树状”缩减操作一样，此缩减代码要求参与共享内存缩减的线程数等于 2 的幂才能正常工作。

请注意，这意味着您可以设计一个缩减内核，它可以通过让最接近的 2 个线程的较小幂执行实际的操作来为每 block 2 个线程的任何 多个 正确运行减少。但是，您发布的代码不能那样工作。