CUDA：如何在 GPU 中将数组的所有元素加总为一个数字？

Question

首先，让我声明我完全知道我的问题已经被问到：Block reduction in CUDA 但是，我希望澄清一下，我的问题是对此的跟进，我有特殊需要使该 OP 找到的解决方案不合适。

那么，让我解释一下。在我当前的代码中，我在 while 循环的每次迭代中运行一个 Cuda 内核，以对数组的值进行一些计算。举个例子，可以这样想：

int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
    __global__ void calcKernel(int* idata, int* odata)
    {
        int i = blockIdx.x*blockDim.x + threadIdx.x;
        if (i < n)
        {
            odata[i] = (idata[i] + 2) * 5;
        }
    }

    iteration++;
}

但是，接下来我必须为 GPU 完成看似艰巨的任务。在调用内核的 while 循环的每次迭代中，我必须对在 odata 中生成的所有值求和，并将结果保存在名为 result 的 intarray 中，该数组中的位置对应于当前迭代。它必须在在内核中或至少仍在GPU中完成，因为由于性能限制，我最终只能检索result数组迭代完成。

一个错误的幼稚尝试看起来像下面这样：

int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
    __global__ void calcKernel(int* idata, int* odata, int* result)
    {
        int i = blockIdx.x*blockDim.x + threadIdx.x;
        if (i < n)
        {
            odata[i] = (idata[i] + 2) * 5;
        }
    }

    result[iteration] = 0;
    for(int j=0; j < max_iterations; j++)
    {
        result[iteration] += odata[j];            
    }

    iteration++;
}

当然，由于 GPU 跨线程分发代码，因此上面的代码不起作用。为了了解如何正确地做到这一点，我一直在网站上阅读有关使用 CUDA 减少阵列的其他问题。特别是，我发现一个非常好的 NVIDIA 的 pdf 关于此类主题的提及，在我开头提到的前一个 SO 问题中也有讨论：http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf

但是，虽然我完全理解此类幻灯片中描述的代码步骤以及一般优化，但我不明白如果代码实际上输出完整，那么该方法如何将数组归约为一个数字数组（以及不明确的维度之一）。有人可以解释一下，并给我一个例子来说明它是如何工作的（即如何从输出数组中取出一个数字）？

现在，回到我在开头提到的那个问题 (Block reduction in CUDA)。请注意，它接受的答案仅建议阅读我上面链接的 pdf - not 谈论如何处理代码生成的输出数组。在 cmets 中，那里的 OP 提到他/她能够通过在 CPU 上对输出数组求和来完成这项工作——这是我无法做到的，因为这意味着我的 while 循环的每次迭代都下载输出数组。最后，该链接中的第三个答案建议使用库来实现这一点 - 但我有兴趣学习这样做的本机方式。

另外，我也会对有关如何实现我上面描述的内容的任何其他提议非常感兴趣。

Answer 1

您已经找到有关块并行减少的规范信息，因此我不会重复。如果您不想自己编写大量新代码来执行此操作，我建议您查看 CUB 库 block_reduce implementation，它提供了一个最佳的块缩减操作，在您现有的代码中添加了大约 4 行代码内核。

关于这里真正的问题，如果你做这样的事情，你可以做你想做的事：

__global__ void kernel(....., int* iter_result, int iter_num) {

    // Your calculations first so that each thread holds its result

    // Block wise reduction so that one thread in each block holds sum of thread results

    // The one thread holding the adds the block result to the global iteration result
    if (threadIdx.x == 0)
        atomicAdd(iter_result + iter_num, block_ressult);
}

这里的关键是atomic function 用于使用给定块的结果安全地更新内核运行结果，而不会出现内存竞争。您绝对必须在运行内核之前初始化iter_result，否则代码将无法运行，但这是基本的内核设计模式。

Answer 2

如果您添加 2 个连续数字并保存结果，在您保存这些数字的任何插槽中，您只需运行多次相同的内核，以不断减少数组总和的 2 次方，例如在这个例子中：

对值求和的数组：

[·1,·2,·3,·4,·5,·6,·7,·8,·9,·10]

首先运行 n/2 个线程，对连续数组元素求和，并将其存储在每个线程的“左侧”，数组现在看起来像：

[·3,2,·7,4,·11,6,·15,8,·19,10]

运行相同的内核，运行 n/4 个线程，现在添加每 2 个元素，并将其存储在最左边的元素中，数组现在看起来像：

[·10,2,7,4,·26,6,15,8,·19,10]

运行同一个内核，运行n/8个线程，现在将每4个元素相加，并存储在数组最左边的元素中，得到：

[·36,2,7,4,26,6,15,8,·19,10]

最后一次运行，单线程将每8个元素相加，并存入数组最左边的元素，得到：

[55,2,7,4,26,6,15,8,19,10]

这样，你只需要以一些线程作为参数来运行你的内核，以获取最后的redux，在第一个元素（55）中查看“点”（·）以查看数组中的哪些元素每次运行都是“活跃的”来总结它们。