OpenCL - 矢量化与线程内循环

我有一个问题，我需要处理已知数量的线程并行（很好），但每个线程可能有大量不同数量的内部迭代（不是很好）。在我看来，这使得它更好地做到这样的核心方案：OpenCL - 矢量化与线程内循环

__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(); 

}

这将必然执行的FULL可能的glIDy的范围内，没有办法提前终止，每这场讨论是：

Killing OpenCL Kernels

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

实际问题：

有没有对付这种在GPU架构比我提出的第一种方案更有效的方法？

我也对此主题的一般信息开放。

谢谢。

来源

2014-01-13 Steve Novakov

假设x尺寸给了你足够的平行度，你在循环变化尺寸的方法应该没问题。请注意，整个内核时间将由最长运行循环决定。如果您可以生成足够的线程，或者如果您有其他可以使用从快速线程组释放的资源开始的并发内核，这可能会很好。 –

我不认为有可以给这个问题一个笼统的回答。这真的取决于你的问题。

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

for循环/如果else语句可能会或可能不会有一个内核的性能产生影响。事实是，性能成本不在内核级别，而是在工作组级别。一个工作组由一个或多个扭曲（NVIDIA）/波阵面（AMD）组成。这些扭曲（我将继续使用NVIDIA的术语，但对于AMD来说则完全一样）都是在锁步执行的。

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

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

现在，知道并根据您的具体问题，您可能能够以这样的方式将您的线程分组，即在工作组内不存在分歧，尽管您在工作组之间存在分歧影响那里）。

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

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

我也喜欢this article。这主要是关于提高性能以简化GPU（不是你的问题）的讨论，但本文的最后一部分还检查了CPU的性能。

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

来源

2014-01-18 13:45:54 CaptainObvious

谢谢，我认为这与我正在寻找的答案非常接近。我想下一代与引入DP并不重要。此外，我觉得需要一个更普遍的PCI-E“APU”卡，例如即将推出的72核心intel之一，但不包括North/Southbridge/SATA等，以及PCIE和DDR3/4接口。这样可以根据问题在GPU，APU和普通CPU之间进行选择。 –

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

两个注意事项：

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

来源

2014-01-13 07:39:49 Oak

那么，我认为这些问题的一般策略，我发现，以及其他人似乎在我看到的各种开源软件中使用的是，有一个主机端线程管理套件，并有效地分解计算成更小的迭代集并检查线程是否在两者之间终止。这方面的问题包括大量的主机设备IO。对我来说，我列出的两个区别是第一个方案可能比第二个方案浪费更少的能量。我诚实地发布了这个完全期待的一堆例子。 –

我在NVIDIA的幻灯片中看到了kepler，有设备内线程排队支持（一个线程可以排队另一个线程），除非我错过了某些东西，这可以有效地解决这个问题。但我认为这不会让它成为AMD的硬件，（或者即将推出），所以在即将到来的OCL版本中看到支持的机会可能很少。 –

http://www.nvidia.com/content/PDF/kepler/NVIDIA-Kepler-GK110-Architecture-Whitepaper.pdf –

OpenCL - 矢量化与线程内循环

回答

相关问题