boosting parallel reduction

Question

我有一个算法，在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];
    }
}

Answer 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 )

它将导致完全顺序访问。