OpenCL：关于 SIMT 执行模型的基本问题答案

【问题标题】：OpenCL: basic questions about SIMT execution modelOpenCL：关于 SIMT 执行模型的基本问题
【发布时间】：2011-09-14 18:12:33
【问题描述】：

“SIMT”架构的一些概念和设计我还不清楚。

据我所见和阅读，完全不同的代码路径和 if() 是一个相当糟糕的主意，因为许多线程可能会以同步方式执行。现在这到底是什么意思？像这样的东西呢：

kernel void foo(..., int flag)
{
    if (flag)
        DO_STUFF
    else
        DO_SOMETHING_ELSE
}

参数“flag”对于所有工作单元都是相同的，并且所有工作单元都采用相同的分支。现在，GPU 是否会执行所有代码，仍然序列化所有内容，并且基本上仍然采用未采用的分支？还是更聪明一点，只要所有线程都同意所采用的分支，就只会执行所采用的分支？这将永远是这种情况。

即序列化总是发生还是仅在需要时发生？对不起这个愚蠢的问题。 ;)

【问题讨论】：

【解决方案1】：

不，并不总是发生。仅当条件在本地工作组中的线程之间不一致时才会执行两个分支，这意味着如果条件在本地工作组中的工作项之间评估为不同的值，则当前一代 GPU 将执行两个分支，但仅执行正确的分支会写入值并有副作用。

因此，保持一致性对于 GPU 分支的性能至关重要。

【讨论】：

我不认为它是至关重要的。它是保持性能的重要因素。虽然这取决于分支的复杂程度，但至少在我目前遇到的大多数情况下，优化内存访问（合并、避免银行冲突、总体上减少流量）比确保分支评估更重要相同的路径（在某些情况下会与这些优化相悖）。
@Grizzly 如果你使用分支是至关重要的 :)

【解决方案2】：

不确定 ati，但对于 nvidia - 它很聪明。如果 warp 中的每个线程都以相同的方式进行，则不会有序列化。

【讨论】：

【解决方案3】：

在您的示例中，flag 对于所有工作项将具有相同的值，因此一个好的编译器将生成代码，将所有工作项带向同一方向。

但考虑以下情况：

kernel void foo(..., int *buffer)
{
    if (buffer[get_global_id(0)])
        DO_STUFF
    else
        DO_SOMETHING_ELSE
}

这里不能保证所有工作项都采用相同的路径，因此需要序列化或控制流消除。

【讨论】：

代码路径是否发散的决定应该根据具体情况在硬件内部进行，这意味着对于您的示例序列化可能会或可能不会发生，具体取决于缓冲区中的值（它可能只某些工作组会发生这种情况，但其他工作组不会）
什么是控制流消除？你能再解释一下吗？谢谢