OpenCL/CUDA：两阶段约简算法

Question

OpenCL/CUDA：两阶段约简算法

3

通过多次调用__reduce()，可以对大数组进行缩减。

然而，以下代码仅使用两个阶段，并在此处文档化:

然而，我无法理解这两个阶段的缩减算法。有人能给一个更简单的解释吗？

__kernel
void reduce(__global float* buffer,
        __local float* scratch,
        __const int length,
        __global float* result) {

    int global_index = get_global_id(0);
    float accumulator = INFINITY;
    // Loop sequentially over chunks of input vector
    while (global_index < length) {
        float element = buffer[global_index];
        accumulator = (accumulator < element) ? accumulator : element;
        global_index += get_global_size(0);
    }

    // Perform parallel reduction
    int local_index = get_local_id(0);
    scratch[local_index] = accumulator;
    barrier(CLK_LOCAL_MEM_FENCE);
    for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
        if (local_index < offset) {
            float other = scratch[local_index + offset];
            float mine = scratch[local_index];
            scratch[local_index] = (mine < other) ? mine : other;
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    if (local_index == 0) {
        result[get_group_id(0)] = scratch[0];
    }
}

使用CUDA也可以很好地实现它。

- gpuguy

1

这个并行约简代码是经典CUDA约简的直接OpenCL移植（缺少一些OpenCL无法执行的优化）。它是由NVIDIA的Mark Harris编写的。您可以在CUDA SDK的约简示例中找到一个非常有教育意义的算法白皮书。阅读完该白皮书后，请编辑您的问题以解释您不理解的内容，然后有人可能会能够提供进一步的帮助。 - talonmies

感谢指出正确的方向。我在理解两阶段和多阶段缩减内核之间的区别时遇到了困难，如此链接所述。 - gpuguy

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Eugene Smith · Accepted Answer

您创建了N个线程。第一个线程查看位置0、N、2N等处的值。第二个线程查看1、N+1、2N+1等处的值。这是第一次循环。它将length个值减少为N个值。

然后，每个线程将最小值保存在共享/本地内存中。然后有同步指令 (barrier(CLK_LOCAL_MEM_FENCE))。接着进行标准的共享/本地内存减少。完成后，本地ID为0的线程将其结果保存在输出数组中。

总而言之，您从length到N/get_local_size(0)值中进行减少。在此代码执行完毕后，您需要进行最后一次传递。但是，这使得大部分工作完成了。例如，您可能具有长度约为10^8，N = 2^16，get_local_size(0) = 256 = 2^8，而此代码将10^8个元素减少为256个元素。

您不理解哪些部分呢？