0

我目前正在开发一个在 NVIDIA Tesla C1060（驱动程序版本 195.17）上起诉 OpenCL 的项目。但是，我遇到了一些我无法真正解释的奇怪行为。这是让我感到困惑的代码（为了清晰和测试目的而减少了）：

kernel void TestKernel(global const int* groupOffsets, global       float* result,     
                       local        int* tmpData,             const int    itemcount)
{
   unsigned int groupid    = get_group_id(0);
   unsigned int globalsize = get_global_size(0);
   unsigned int groupcount = get_num_groups(0);

   for(unsigned int id = get_global_id(0); id < itemcount; id += globalsize, groupid += groupcount)
   {
      barrier(CLK_LOCAL_MEM_FENCE);
      if(get_local_id(0) == 0)
         tmpData[0] = groupOffsets[groupid]; 
      barrier(CLK_LOCAL_MEM_FENCE);
      int offset = tmpData[0];
      result[id]   = (float) offset;
   }
}

此代码应将每个工作组的偏移量加载到本地内存中，然后将其读回并将其写入相应的输出向量条目。对于大多数工作项，这是可行的，但对于每个工作组，本地 ID 为 1 到 31 的工作项读取的值不正确。我的输出向量（对于 workgroupsize=128）如下：

index       0: 0
index   1- 31: 470400
index  32-127: 0
index     128: 640
index 129-159: 471040
index 160-255: 640
index     256: 1280
index 257-287: 471680
index 288-511: 1280
...

我期望的输出是

index   0-127: 0
index 128-255: 640
index 256-511: 1280
...

奇怪的是：只有当我使用少于 itemcount 的工作项时才会出现问题（因此当 globalsize>=itemcount 时它按预期工作，这意味着每个工作项只处理一个条目）。所以我猜它与循环有关。有谁知道我做错了什么以及如何解决？

更新：我发现如果我改变它似乎工作

if(get_local_id(0) == 0)
     tmpData[0] = groupOffsets[groupid];

至

if(get_local_id(0) < 32)
     tmpData[0] = groupOffsets[groupid];

这更让我吃惊，所以虽然它可能会解决问题，但我觉得以这种方式修复它并不自在（因为它可能会在其他时候中断）。此外，我宁愿避免在 Geforce 8xxx 类硬件上运行时由于额外的（据我了解的硬件未合并）内存访问而损失性能。所以问题仍然存在。

4

1 回答 1

1

首先，重要的是，您需要小心它itemcount是本地工作大小的倍数，以避免在执行屏障时出现分歧。

在处理器上执行内核的工作组中的所有工作项必须执行此功能，然后才能允许任何工作项继续执行越过障碍。执行内核的工作组中的所有工作项都必须遇到此函数。

您可以按如下方式实现：

unsigned int itemcountrounded = get_local_size(0) * ((itemcount + get_local_size(0) - 1) / get_local_size(0));
for(unsigned int id = get_global_id(0); id < itemcountrounded; id += globalsize, groupid += groupcount)
{
    // ...
    if (id < itemcount)
        result[id]   = (float) offset;
}

你说为了简单起见减少了代码，如果你运行你发布的内容会发生什么？只是想知道您是否也需要在全局内存上设置障碍。

于 2010-02-01T20:49:21.740 回答