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 的 FULL POSSIBLE RANGE 范围内执行，而无法提前终止：

Killing OpenCL Kernels

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

实际问题：

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

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

谢谢。

Answer 1

我认为这个问题无法给出一个普遍的答案。这真的取决于你的问题。

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

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

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

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

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

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

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

我也喜欢this article。它主要是关于提高性能以简单减少 GPU 的讨论（不是您的问题），但文章的最后部分还检查了 CPU 上的性能。

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

Answer 2

我更喜欢第二个版本，因为for 在迭代之间插入了错误的依赖关系。如果内部迭代是独立的，则将每个迭代发送到不同的工作项，并让 OpenCL 实现解决如何最好地运行它们。

两个警告：

• 如果平均迭代次数明显低于最大迭代次数，这可能不值得额外的虚拟工作项。
• 您将有更多的工作项，您仍然需要计算每个工作项的条件...如果计算条件很复杂，这可能不是一个好主意。
  • 或者，您可以将索引展平到 x 维度，将所有迭代分组到同一个工作组中，然后每个工作组只计算一次条件，并使用本地内存 + 障碍来同步它。