CUDA条件线程同步答案

【问题标题】：CUDA conditional thread synchronizationCUDA条件线程同步
【发布时间】：2011-07-23 20:51:38
【问题描述】：

CUDA 编程指南指出

__syncthreads() 在条件代码中是允许的，但前提是条件评估相同跨越整个线程块，否则代码执行很可能挂起或产生意外的一面效果。

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

我想可能有各种情况下您可能需要这样做；例如，如果您有一个二进制掩码并且需要有条件地对像素应用某个操作。说，if (mask(x, y) != 0) 然后执行包含__syncthreads() 的代码，否则什么也不做。那要怎么做呢？

【问题讨论】：

标签： c++ parallel-processing cuda synchronization

【解决方案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

【讨论】：

这个答案有点过时了，但它在搜索结果中的排名很高。我发现另一个答案与今天的 CUDA 功能更相关：stackoverflow.com/questions/15146886/…。也许这个答案需要更新？
@Liang：理论上这个答案还是正确的； CUDA 模型规定块内的所有线程都必须到达屏障。您引用的帖子描述了为什么提前退出有效，但它不适用于所有 GPU (G80)，并且不能保证它总是有效。

【解决方案2】：

从 3.0 开始，您可以使用 warp vote 函数来完成 __syncthreads 无法完成的任务：

仅支持 Warp 投票功能按计算能力 1.2 的设备

int __all(int predicate); 谓词对于经线的所有线程和当且仅当返回非零谓词计算为非零全部。

int __any(int 谓词); 评估谓词经线和返回的所有线程非零当且仅当谓词其中任何一个的计算结果都为非零。

unsigned int __ballot(int predicate); 评估所有线程的谓词扭曲并返回一个整数，其当且仅当第 N 位被设置谓词计算为非零经线的第 N 根线。这功能仅受设备支持计算能力 2.x.

否则还有 Atomic Bitwise 函数

atomicAnd、atomicOr、atomicXor

参见 cuda 编程指南的 B.11 部分

【讨论】：