加速并行减少OpenCL

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

谢谢！

__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]; 
    } 
}

来源

2013-04-27 Valera

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

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

global_index += get_global_size(0);

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

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

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

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

这将导致完全顺序访问。

来源

2013-04-28 12:14:42

加速并行减少OpenCL

回答

相关问题