在 CUDA 中找到最大值答案

【问题标题】：Finding max value in CUDA在 CUDA 中找到最大值
【发布时间】：2011-07-12 11:23:34
【问题描述】：

我正在尝试在 CUDA 中编写代码来查找最大值对于给定的一组数字。

假设您有 20 个数字，并且内核在 2 个 5 个线程的块上运行。现在假设 10 个线程同时比较前 10 个值，线程 2 找到一个最大值，因此线程 2 正在更新全局内存中的最大值变量。在线程 2 更新时，将使用旧值进行比较的剩余线程 (1,3-10) 会发生什么情况？

如果我使用 atomicCAS() 锁定全局变量，线程 (1,3-10) 是否会使用旧的最大值进行比较？我该如何克服这个问题？

【问题讨论】：

标签： parallel-processing cuda reduction

【解决方案1】：

这纯粹是一个归约问题。这是 NVIDIA 提供的一个很好的 presentation，用于优化 GPU 上的减少。您可以使用相同的技术来找到所有元素的最小值、最大值或总和。

【讨论】：

与演示文稿相关的问题：Sum reduction with CUDA: What is N? 接受的答案表明最终内核代码中有错字。

【解决方案2】：

Thrust 库的链接已损坏。
如果有人觉得在这种情况下使用它很有用，您可以在此处找到文档：
Thrust, extrema reductions

【讨论】：

【解决方案3】：

除非您尝试编写缩减内核，否则最简单的方法就是使用CUBLAS。

【讨论】：

【解决方案4】：

我寻找了相同的答案，但发现大多数对于像我这样的新手来说都太强大了。这是我查找最大值的示例代码。请让我知道这是否正确使用。

__global__
void find_max(int max_x, int max_y, float *tot, float *x, float *y)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    int j = blockIdx.y*blockDim.y + threadIdx.y;
    if(i < max_x && j<max_y) {
        if(*tot < x[i])
            atomicExch(tot, x[i]);
    }
}

【讨论】：

那里可能存在竞争条件：比较成功，另一个线程存储一个新的最大值，然后这个线程用它的最大值覆盖它（可能小于另一个线程存储的）。这就是为什么问题提到CAS，而不是Exch。（此外，您无缘无故地修改了x[i]，并且您的函数甚至不使用y[] 参数。）将所有内容原子地交换到一个共享计数器中而不是在多个点上设置局部最大值，这看起来也非常低效然后在最后结合来自多个线程的结果。（我不知道 CUDA 所以 IDK 怎么做。）