OpenCL：SIMT执行模型的基本问题

Question

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

根据我所看到和阅读的内容，不同的代码路径和 if() 一起是一个相当糟糕的主意，因为许多线程可能会同步执行。那么这究竟意味着什么呢？怎么样：

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

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

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

Answer 1

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

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

Answer 2

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

Answer 3

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

但请考虑以下情况：

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

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