algorithm - 在 OpenCL 中查找 float 组最大值的快速方法

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() 之后，只需让运行在 thread_id == 0 上的内核循环遍历减少的结果并找到其中的最大值。这不是最佳解决方案，但可能适合您当前的内核。