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

Posted 2023-04-15

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

【英文标题】：CUDA: how to sum all elements of an array into one number within the GPU?

【发布时间】：2017-07-20 10:12:24

【问题描述】：

首先，让我声明我完全知道我的问题已经被问到：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 循环的每次迭代都下载输出数组。最后，该链接中的第三个答案建议使用库来实现这一点 - 但我有兴趣学习这样做的本机方式。

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

【问题讨论】：

您会考虑使用多个内核而不是单个内核吗？

@RicardoOrtegaMagaña 当然，只要我不从/向 GPU 内存传输，调用多个内核本身不会有问题。

我建议使用 2 个内核，1 个用于第一次计算，另一个内核仅用于在主机程序中添加和放置循环。我需要有关您的程序如何工作的更多信息，但是使用您向我们展示的代码，这可能是解决您的问题的一种简单方法。

@RicardoOrtegaMagaña 我现在正在尝试这样做，但我在最后一步中挣扎。与我链接到的问题中最初提出的问题完全相同（但在 CPU 端已解决）。是的，我现在使用内核来进行缩减，但最终得到的是一个数组，而不是一个数字，我想在 GPU 中进一步将这个最终数组缩减为一个数字。

在我给你的例子中，结果它在数组的第一个元素上，如果你想在单个变量（内存位置）中发送一个指针在内核中，并将该值分配给那个位置，你想要那个值的地方，这样，当你想从主机或设备读取它时，你只需给出那个指针。

【参考方案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，否则代码将无法运行，但这是基本的内核设计模式。

【讨论】：

非常有见地。我只是能够正确地实现它并且它有效。就我而言，唯一的缺点是我正在使用字符，而 atomicAdd 不适用于字符，所以我必须转换相关部分，浪费一点内存和一点性能。尽管如此，在“atomicAdd”与“仅在输出数组中保存块结果”的纯粹比较中，我对使用 atomicAdd 所损失的性能很少感到惊讶。

【参考方案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）中查看“点”（·）以查看数组中的哪些元素每次运行都是“活跃的”来总结它们。

【讨论】：

感谢您的回答。虽然我理解逻辑，但我不知道如何实现它。你能给我举个小例子吗？看起来您所说的与pdf我在问题中链接的内容完全相同。