CUDA条件线程同步

Question

CUDA编程指南指出：

__syncthreads()在条件代码中是允许的，但只有当条件在整个线程块中的评估是相同的时才可以，否则代码执行可能会挂起或产生意外的副作用。

因此，如果我需要通过一个块进行条件分支来同步线程，其中一些线程可能会或可能不会采取包括__syncthreads()调用的分支，这是否意味着它无法工作？

我想象中可能会有各种情况需要这样做；例如，如果您有一个二进制掩码，并且需要有条件地对像素应用某个操作。比如，如果(mask(x, y) != 0)，则执行包含__syncthreads()的代码，否则不执行任何操作。该如何实现？

- Kristian D'Amato

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

- Tom

---

这个答案有点过时，但它在搜索结果中排名很高。我发现另一个答案在今天的CUDA功能中更具相关性：https://dev59.com/H2_Xa4cB1Zd3GeqPwxGQ 。也许这个答案需要更新一下？ - Liang

@Liang：从理论上讲，这个答案仍然是正确的；CUDA模型规定块内所有线程必须到达屏障。你引用的帖子解释了为什么提前退出有效，但它并不适用于所有GPU（G80），也不能保证它总是有效。 - Tom

Answer 2

从3.0版本开始，您可以使用Warp投票函数来完成__syncthreads无法实现的功能：

仅支持计算能力为1.2的设备

int __all(int predicate); 谓词 应用于warp的所有线程，并且如果且仅当谓词对于它们 所有的线程都求值为非零时返回非零。

int __any(int predicate); 应用于warp的所有线程并返回 非零如果且仅当谓词对于任何一个 线程求值为非零。

unsigned int __ballot(int predicate); 应用于warp的所有线程并返回一个整数， Nth比特设置为仅当谓词对于第N个 线程求值为非零。此 函数仅受支持 的设备具有计算能力2.x。

否则也有原子位运算函数

atomicAnd、atomicOr、atomicXor

请参阅cuda编程指南的B.11节

- fabrizioM