OpenCL的浮子总和减少

我想申请在这片我的内核的代码减少（1个维数据）：OpenCL的浮子总和减少

__local float sum = 0; 
int i; 
for(i = 0; i < length; i++) 
    sum += //some operation depending on i here;

代替具有正好1个线程执行此操作的，我想有n个线程（n =长度），最后有1个线程作总和。

在伪代码，我想能写这样的事：

int i = get_global_id(0); 
__local float sum = 0; 
sum += //some operation depending on i here; 
barrier(CLK_LOCAL_MEM_FENCE); 
if(i == 0) 
    res = sum;

有没有办法？

我总有一个竞赛条件。

来源

2013-12-16 Kami

这就是所谓的并行减少，查找它。它不像你的片段那么容易，但也不是很难。只需要更多的工作。 – Thomas

为了让你开始你可以做下面的例子（see Scarpino）。这里我们还利用了OpenCL float4数据类型的矢量处理。

请记住，下面的内核返回一些部分和：每个本地工作组返回主机。这意味着您必须将所有部分总和重新计算在主机上，才能执行最终总和。这是因为（至少在OpenCL 1.2中）没有障碍函数来同步不同工作组中的工作项。

如果总结主机上的部分总和是不可取的，那么可以通过启动多个内核来解决此问题。这会引入一些内核调用开销，但在某些应用程序中，额外的惩罚是可接受的或不重要的。要用下面的例子来做到这一点，您需要修改主机代码以重复调用内核，然后在输出向量数量低于本地大小后停止执行内核（详细信息请咨询您或查看Scarpino reference）。

编辑：为输出添加额外的内核参数。添加了点积来总结浮点数4向量。

__kernel void reduction_vector(__global float4* data,__local float4* partial_sums, __global float* output) 
{ 
    int lid = get_local_id(0); 
    int group_size = get_local_size(0); 
    partial_sums[lid] = data[get_global_id(0)]; 
    barrier(CLK_LOCAL_MEM_FENCE); 

    for(int i = group_size/2; i>0; i >>= 1) { 
     if(lid < i) { 
      partial_sums[lid] += partial_sums[lid + i]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
    } 

    if(lid == 0) { 
     output[get_group_id(0)] = dot(partial_sums[0], (float4)(1.0f)); 
    } 
}

来源

2013-12-16 15:03:11

减少数据的简单而快速的方法是将数据的上半部分反复折叠到下半部分。

例如，请使用以下可笑的简单CL代码：

__kernel void foldKernel(__global float *arVal, int offset) { 
    int gid = get_global_id(0); 
    arVal[gid] = arVal[gid]+arVal[gid+offset]; 
}

随着下面的Java/JOCL主机代码（或它的端口，以C++等）：

int t = totalDataSize; 
    while (t > 1) { 
     int m = t/2; 
     int n = (t + 1)/2; 
     clSetKernelArg(kernelFold, 0, Sizeof.cl_mem, Pointer.to(arVal)); 
     clSetKernelArg(kernelFold, 1, Sizeof.cl_int, Pointer.to(new int[]{n})); 
     cl_event evFold = new cl_event(); 
     clEnqueueNDRangeKernel(commandQueue, kernelFold, 1, null, new long[]{m}, null, 0, null, evFold); 
     clWaitForEvents(1, new cl_event[]{evFold}); 
     t = n; 
    }

主机代码循环log2（n）次，因此即使使用巨大的数组也会快速完成。带“m”和“n”的小提琴是用来处理非幂次数组的。

易于OpenCL平行良好的任何GPU平台（即快速）。
低内存，因为它
工程有效地工作在适当位置与非幂的两个数据大小
柔性的，例如您可以更改内核以执行“min”而不是“+”

来源

2018-01-02 09:52:54

如果您支持OpenCL C 2，则可以使用新的work_group_reduce_add()函数减少单个工作组内的总和。0功能

来源

2018-02-22 09:33:26

OpenCL的浮子总和减少

回答

相关问题