2

我正在使用 OpenCL。我对以下示例中如何执行工作项感兴趣。

我的一维范围为 10000，工作组大小为 512。内核如下：

__kernel void
doStreaming() {
  unsigned int id = get_global_id(0);

  if (!isExecutable(id))
    return;

 /* do some work */
}

在这里它检查是否需要使用以下 id 处理元素。

假设执行从 512 大小的第一个工作组开始，其中 20 个被拒绝isExecutable。GPU 是否继续执行其他 20 个元素而不等待前 492 个元素？

不涉及任何障碍或其他同步技术。

4

2 回答 2

2

当一些工作项偏离通常的 /* 做一些工作 */ 时，它们可以通过从下一个波前（amd）或下一个经线（nvidia）获取指令来利用管道占用优势，因为当前的经线/波前工作项正忙于做其他事情。但这会导致内存访问序列化并清除工作组的访问顺序，从而降低性能。

避免出现发散的扭曲/波前：如果你在循环中执行 if 语句，那真的很糟糕，所以你最好找到另一种方法。

如果工作组中的每个工作项都具有相同的分支，则可以。

如果每个工作项在数百次计算中只进行很少的分支，那没关系。

尝试为所有工作项（令人难以置信的并行数据/算法）生成相同的条件，以利用 gpu 所拥有的能力。

我知道摆脱最简单的分支与计算情况的最佳方法是使用全局是-否数组。0=yes, 1=no ：总是计算，然后将结果与工作项的 yes-no 元素相乘。通常，每个核心添加 1 字节元素内存访问要比每个核心执行一个分支要好得多。实际上，在添加这个 1 字节后，将对象长度设为 2 的幂可能会更好。

于 2013-07-07T22:57:41.097 回答
1

是和不是。以下详细说明基于 NVIDIA 的文档，但我怀疑它在 ATI 硬件上是否有任何不同（尽管实际数字可能会有所不同）。通常，工作组的线程在所谓的 warp 中执行，是工作组大小的子块。在 NVIDIA 硬件上，每个工作组被分成每个 32 个线程的 warp。并且这些扭曲中的每一个都以锁步执行，因此完全并行（它可能不是实时并行的，这意味着可能有 16 个并行线程，然后直接再次出现 16 个线程，但从概念上讲，它们完全并行运行） . 因此，如果这 32 个线程中只有一个执行该附加代码，则其他线程将等待它。但是所有其他经线中的线程不会关心所有这些。

所以是的，可能有线程会不必要地等待其他线程，但这发生的规模小于整个工作组的规模（任何 NVIDIA 硬件上为 32 个）。这就是为什么如果可能的话应该避免内部warp分支偏差的原因，这也是为什么保证在单个warp中工作的代码不需要任何同步，例如共享内存访问（算法的常见优化）。

于 2013-07-07T23:04:20.737 回答