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的浮子總和減少

回答

相關問題