提升并行缩减 OpenCL答案

【问题标题】：boosting parallel reduction OpenCL提升并行缩减 OpenCL
【发布时间】：2013-04-28 20:14:42
【问题描述】：

我有一个算法，在 GPU 上执行两阶段并行缩减以找到字符串中的最小元素。我知道有一个关于如何让它工作得更快的提示，但我不知道它是什么。关于如何调整这个内核以加快我的程序的任何想法？不需要实际更改算法，可能还有其他技巧。欢迎所有想法。

谢谢！

__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
        while (global_index < length) {
            float element = buffer[global_index];
            accumulator = (accumulator < element) ? accumulator : element;
            global_index += get_global_size(0);
    }
    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];
    }
}

【问题讨论】：

标签： opencl gpu

【解决方案1】：

accumulator = (accumulator < element) ? accumulator : element;

使用fmin 函数 - 这正是您所需要的，它可能会导致更快的代码（调用内置指令，如果可用，而不是昂贵的分支）

global_index += get_global_size(0);

你的典型get_global_size(0)是什么？

虽然您的访问模式不是很糟糕（它是合并的，32-warp 的 128 字节块） - 最好尽可能按顺序访问内存。例如，顺序访问可能有助于memory prefetching（注意，OpenCL 代码可以在任何设备上执行，包括 CPU）。

考虑以下方案：每个线程将处理范围

[ get_global_id(0)*delta ,  (get_global_id(0)+1)*delta )

这将导致完全顺序访问。

【讨论】：