1

很长一段时间以来，我一直在尝试进行简单的扫描。对于小问题，输出是正确的，但是对于大输出，我有时会得到正确的结果。我检查了Apple 的 OpenCL 示例，我基本上在做同样的事情（除了银行冲突，我忽略了 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];
    }
  }
}

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

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

我的猜测是，不同线程同时访问相同数据是一个问题，但是，情况并非如此，因为每个工作组都在处理不同的输入数据块。对此事的任何帮助将不胜感激！

4

1 回答 1

4

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

您的块变量是 get_group_id(0)/2，这意味着您有两个完整的工作组分配给同一个块。你可能想要反过来。如果它们碰巧同步运行，它们将简单地覆盖彼此的工作，因为它们的加载存储依赖关系将匹配。否则，它们可能会或可能不会干扰，总是在多次求和的方向上。

你的问题本身就有一个提示：“每个工作组都可以扫描两倍大小的块。” 这应该意味着数组大小一半的总工作大小就足够了。

downsweep() 中的循环也有一个奇怪的地方。第一次迭代什么都不做；lid+1>=1，并且 d 从 1 开始。这可能是一个微不足道的多余迭代，但它在计划中是一个偏差。

于 2010-07-23T02:43:28.063 回答