OpenCL/CUDA：两阶段归约算法

Question

可以通过调用 __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 很好地实现。

Answer 1

您创建N线程。第一个线程查看位置 0, N, 2*N, ... 第二个线程查看值 1, N+1, 2*N+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 个元素.

哪些部分你不明白？