c++ - OpenCL - 向量化与线程内 for 循环

Question

我遇到一个问题，我需要并行处理已知数量的线程(很好)，但每个线程的内部迭代次数可能大相径庭(不太好)。在我看来，这样的内核方案会更好:

__kernel something(whatever)
{
   unsigned int glIDx = get_global_id(0);

   for(condition_from_whatever)
   {

   }//alternatively, do while

}

其中 id(0) 是事先已知的，而不是:

__kernel something(whatever)
{
   unsigned int glIDx = get_global_id(0);
   unsigned int glIDy = get_global_id(1); // max "unroll dimension"

   if( glIDy_meets_condition)
      do_something();
   else
      dont_do_anything();

}

根据这个讨论，它必须在 glIDy 的全部可能范围内执行，没有办法提前终止:

Killing OpenCL Kernels

我似乎找不到任何关于内核中动态大小的 forloops/do-while 语句成本的具体信息，尽管我确实在 Nvidia 和 AMD 的 SDK 的内核中随处看到它们。我记得读过一些关于内核内条件分支越不周期性，性能越差的文章。

实际问题:

有没有比我提出的第一个方案更有效的方法在 GPU 架构上处理这个问题？

我也愿意接受有关此主题的一般信息。

谢谢。

Answer 1

我认为这个问题没有一个通用的答案。这实际上取决于您的问题。

但是这里有一些关于这个主题的注意事项:

for 循环/if else 语句可能会或可能不会影响内核的性能。事实上，性能成本不在内核级别，而是在工作组级别。一个工作组由一个或多个 warp (NVIDIA)/wavefront (AMD) 组成。这些扭曲(我将保留 NVIDIA 术语，但它与 AMD 完全相同)以锁步方式执行。

因此，如果在 warp 中由于 if else(或具有不同迭代次数的 for 循环)而出现分歧，则执行将被序列化。也就是说，这个 warp 中的线程遵循第一条路径将完成它们的工作，而其他线程将空闲。一旦他们的工作完成，这些线程将空闲，而其他线程将开始工作。

如果您需要将线程与屏障同步，则这些语句会产生另一个问题。如果不是所有线程都遇到障碍，您将有未定义的行为。

现在，知道这一点并根据您的具体问题，您可能能够以工作组内没有分歧的方式对线程进行分组，尽管工作组之间会有分歧(没有影响那里)。

还知道 warp 由 32 个线程和 64 个波前组成(可能不是在旧的 AMD GPU 上 - 不确定)，您可以使组织良好的工作组的规模等于或这些数字的倍数。请注意，它非常简单，因为还应考虑其他一些问题。参见例如 this question以及 Chanakya.sun 给出的答案(也许对该主题进行更多挖掘会很好)。

如果您的问题无法按照刚才描述的那样进行组织，我建议考虑在 CPU 上使用 OpenCL，它可以很好地处理分支。如果我没记错的话，通常每个工作组都有一个工作项。在这种情况下，最好查看 Intel 和 AMD 的 CPU 文档。我也很喜欢Heterogeneous Computing with OpenCL的第六章这解释了在编程时将 OCL 与 GPU 和 CPU 结合使用的区别。

我喜欢this article也。它主要是关于通过简单地减少 GPU 来提高性能的讨论(不是你的问题)，但本文的最后一部分还检查了 CPU 的性能。

最后一件事，关于您对@Oak 提供的关于“设备内线程队列支持”的回答的评论，这实际上称为动态并行性。此功能显然可以解决您的问题，但即使使用 CUDA，您也需要具有 3.5 或更高功能的设备。因此，即使是具有 Kepler GK104 架构的 NVIDIA GPU 也不支持它(功能 3.0)。对于 OCL，动态并行是标准版本 2.0 的一部分。 (据我所知还没有实现)。