opencl - OpenCL：关于 SIMT 执行模型的基本问题

Question

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

从我所看到和阅读的内容来看，完全不同的代码路径和 if() 是一个相当糟糕的主意，因为许多线程可能会步调一致地执行。现在这到底是什么意思？怎么样的东西：

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

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

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

score 3 · Accepted Answer

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

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

score 1 · Accepted Answer

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

score 1 · Accepted Answer

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

但考虑以下情况：

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

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

3 回答 3