在OpenCL中执行扫描

Question

我一直试图让一个简单的扫描工作很长一段时间。对于小问题，输出是正确的，但是对于大输出，我有时只能得到正确的结果。我检查了Apple's OpenCL example，我基本上做同样的事情（除了银行冲突，我忽略了atm）。因此，这里的第一个阶段的代码：

__kernel void 
scan_init(__global int * input, 
      __global int * sums) 
{ 
    int gid = get_global_id(0); 
    int lid = get_local_id(0); 
    int chunk_size = get_local_size(0)*2; 

    int chunk = gid/chunk_size; 
    int offset = chunk*chunk_size; 

    reduction(input, offset); 

    // store sums 
    if(lid==0) 
    { 
    sums[chunk] = input[(chunk+1)*chunk_size-1]; 
    } 

    downsweep(input, offset); 
} 

，还原功能本身：

void reduction(__global int * input, 
     int offset) 
{ 
int stride = 1; 
int grp_size = get_local_size(0); 
int lid = get_local_id(0); 

for(int d = grp_size; d > 0; d>>=1) 
{ 
    barrier(CLK_GLOBAL_MEM_FENCE); 

    if(lid < d) 
    { 
    int ai = stride*(2*lid+1)-1+offset; 
    int bi = stride*(2*lid+2)-1+offset; 
    input[bi] += input[ai]; 
    } 

    stride *= 2; 
    } 
} 

在第二阶段，泛音资金用于构建总和为每个元素：

void downsweep(__global int * input, 
     const unsigned int offset) 
{ 
    int grp_size = get_local_size(0); 
    int lid = get_local_id(0); 
    int stride = grp_size*2; 

    for(int d = 1; d <= grp_size; d *=2) 
    { 
    barrier(CLK_GLOBAL_MEM_FENCE); 

    stride >>=1; 

    if(lid+1 < d) 
    { 
     int src = 2*(lid + 1)*stride-1+offset; 
     int dest = src + stride; 
     input[dest]+=input[src]; 
    } 
    } 
} 

输入被填充为本地工作量的倍数。每个工作组都可以扫描两倍大小的块。我将每个块的总和保存在总和数组中，我用它来检查结果。以下是1的阵列的输出输入尺寸4000：

Chunk size: 1024 
Chunks: 4 
Scan global size: 4096 
Local work size: 512 
Sum size: 4 
0:1024 1:1120 2:2904 3:928 

然而，预期的结果将是

0:1024 1:1024 2:1024 3:928 

如果我再次运行该代码，我得到：

0:1056 1:5376 2:1024 3:928 
0:1024 1:1088 2:1280 3:992 
0:5944 1:11156 2:3662 3:1900 
0:7872 1:1056 2:2111 3:1248 

调用内核如下：

clEnqueueNDRangeKernel(cl_ctx->queue, scan_init, 1, NULL, &scan_global_size, &local_work_size, 0, NULL, NULL); 

如果全球规模是4096和当地大小为512。如果我限制了本地工作组大小为64，输出如下：

0:128 1:128 2:128 3:288 4:128 5:128 6:192 7:192 
8:192 9:254 10:128 11:256 12:128 13:360 14:128 15:128 
16:128 17:128 18:128 19:288 20:128 21:128 22:128 23:128 
24:192 25:128 26:128 27:192 28:128 29:128 30:128 31:32 

如果我改变输入尺寸为512和任何块大小，一切都很好！最后，当使用输入大小513和256的组大小（也就是说，我有两个块，每个块都有512个元素，第二个块只有第一个元素设置为1）时，第一个元素的结果阶段是：

0:1 1:2 2:1 3:6 4:1 5:2 6:1 7:14 
8:1 9:2 10:1 11:6 12:1 13:2 14:1 15:28 
16:1 17:2 18:1 19:6 20:1 21:2 22:1 23:14 
24:1 25:2 26:1 27:6 28:1 29:2 30:1 31:56 
32:1 33:2 34:1 35:6 36:1 37:2 38:1 39:14 
40:1 41:2 42:1 43:6 44:1 45:2 46:1 47:28 
48:1 49:2 50:1 51:6 52:1 53:2 54:1 55:14 
56:1 57:2 58:1 59:6 60:1 61:2 62:1 63:148 

，它应该是：

0:1 1:2 2:1 3:4 4:1 5:2 6:1 7:8 
8:1 9:2 10:1 11:4 12:1 13:2 14:1 15:16 
16:1 17:2 18:1 19:4 20:1 21:2 22:1 23:8 
24:1 25:2 26:1 27:4 28:1 29:2 30:1 31:32 
32:1 33:2 34:1 35:4 36:1 37:2 38:1 39:8 
40:1 41:2 42:1 43:4 44:1 45:2 46:1 47:16 
48:1 49:2 50:1 51:4 52:1 53:2 54:1 55:8 
56:1 57:2 58:1 59:4 60:1 61:2 62:1 63:64 

我的猜测是，它是由不同的线程同时访问相同数据的问题，但是，这不应该是这样，因为每个工作组正在处理不同的输入数据块。任何关于此事的帮助将不胜感激！

Answer 1

我怀疑问题与barrier（）不是一个工作组间同步。每个工作组都有自己的障碍，而且您对工作组本身的排序没有任何保证。当您将输入集大小更改为512时，您可能会让您的所有工作组在同一个多处理器上运行，并因此被偶然同步。

你的块变量是get_group_id（0）/ 2，这意味着你有两个完整的工作组分配给同一个块。你可能想要另一种方式。如果它们碰巧以锁步方式运行，它们会简单地覆盖对方的工作，因为它们的加载存储库依赖关系会匹配。否则，它们可能会或可能不会干扰，总是在多次求和值的方向上进行干扰。

在这个问题上的一个提示是在你自己的问题中：“每个工作组都可以扫描两倍大小的块。”这应该意味着数组大小的一半是足够的。

下降扫描（）中的循环也具有奇特性。第一次迭代什么也不做;盖+ 1> = 1，并且d从1开始。这可能是一个无足轻重的多余迭代，但是在规划中它是一个偏差。