OpenCL-循环的矢量化与线程内

OpenCL - Vectorization vs In-thread for loop

本文关键字：线程 矢量化 循环 OpenCL-      更新时间：2023-10-16

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

__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的完全可能范围内执行，没有办法提前终止：

杀死OpenCL内核

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

实际问题：

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

我也对这个主题的一般信息持开放态度。

谢谢。

我认为这个问题没有一个通用的答案。这真的取决于你的问题。

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

for循环/if-else语句可能会也可能不会对内核的性能产生影响。事实上，性能成本不是在内核级别，而是在工作组级别。工作组由一个或多个翘曲(NVIDIA)/波前(AMD)组成。这些扭曲(我将保留NVIDIA的术语，但对AMD来说完全相同)是以锁定步骤执行的。

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

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

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

还知道一个翘曲是由32个线程和64个波前组成的(可能不是在旧的AMD GPU上——不确定)，你可以使组织良好的工作组的大小等于或是这些数字的倍数。请注意，它被简化了，因为应该考虑其他一些问题。例如，看看这个问题和Chanakya.sun给出的答案(也许更多地挖掘这个话题会很好)。

如果你的问题不能像刚才描述的那样组织起来，我建议你考虑在处理分支的CPU上使用OpenCL。如果我还记得的话，通常每个工作组都有一个工作项。在这种情况下，最好查看英特尔和AMD的CPU文档。我也非常喜欢OpenCL的异构计算的第6章，它解释了在编程时将OCL与GPU和CPU一起使用的区别。

我也喜欢这篇文章。这主要是关于通过简单地降低GPU来提高性能的讨论(不是你的问题)，但文章的最后一部分也考察了CPU的性能。

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

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

两个注意事项：

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