在 OpenCL 中查找浮点数组最大值的快速方法

Question

我在 OpenCL 中查找数组最大值的简单任务时遇到了麻烦。

__kernel void ndft(/* lots of stuff*/)
{
    size_t thread_id = get_global_id(0); // thread_id = [0 .. spectrum_size[
    /* MATH MAGIC */
    // Now I have    float spectrum_abs[spectrum_size]    and
    // I want the maximum as well as the index holding the maximum
    barrier();
    // this is the old, sequential code:
    if (*current_max_value < spectrum_abs[i])
    {
        *current_max_value = spectrum_abs[i];
        *current_max_freq = i;
    }
}

现在我可以像在单核系统上一样添加if (thread_id == 0)并遍历整个事情，但由于性能是一个关键问题（否则我不会在 GPU 上进行频谱计算），有没有更快的方法做到这一点？

在

上面的内核末尾返回到 CPU 不是一个选项，因为内核实际上在那之后继续。

Answer 1

您需要编写一个并行缩减。将"大"数组拆分为小块（单个工作组可以有效处理的大小），并计算每个数组的最小值。

以

迭代方式执行此操作（涉及主机和设备代码），直到只剩下一组最小/最大值。

请注意，您可能需要编写一个单独的内核来执行此操作，除非当前的工作分发适用于内核的这一部分（请参阅我上面向您提出的问题）。

如果您当前的工作分布是合适的，另一种方法是在每个工作组内找到最小最大值，并将其写入全局内存中的缓冲区（索引 = local_id）。在 barrier（） 之后，只需让内核在 == 0 上运行thread_id遍历减少的结果并找到其中的最大值。这不是最佳解决方案，但可能是适合当前内核的解决方案。