CUDA条件线程同步

Question

CUDA编程指南指出了这一点

条件代码中允许__syncthreads(),但仅当条件在整个线程块中进行相同的求值时,否则代码执行可能会挂起或产生意外的副作用.

因此,如果我需要在一个块上使用条件分支同步线程,其中一些线程可能会或可能不会采用包含该__syncthreads()调用的分支,这是否意味着它不起作用？

我想象可能存在各种各样的情况,你可能需要这样做; 例如,如果您有二进制掩码并需要有条件地对像素应用某个操作.比如说,if (mask(x, y) != 0)然后执行包含的代码__syncthreads(),否则什么都不做.怎么办？

Answer 1

如果你需要走这条路,你可以把身体分成两个阶段:

if (condition)
{
    // code before sync
}
__syncthreads();
if (condition) // or remember a flag or whatever
{
    // code after sync
}

或者,您可以使用条件设置禁用某些操作的标志,例如,如果您正在计算增量更新,则可以执行以下操作:

// *ALL* compute a delta update, those threads that would have failed the condition
// simply compute garbage.
// This can include syncthreads
if (condition)
    // apply update