CUDA：同步线程

几乎在我读到的有关 CUDA 编程的任何地方都提到了 warp 中的所有线程都执行相同操作的重要性。
在我的代码中，我遇到了无法避免某种条件的情况。它看起来像这样：

// some math code, calculating d1, d2
if (d1 < 0.5)
{
    buffer[x1] += 1;  // buffer is in the global memory
}
if (d2 < 0.5)
{
    buffer[x2] += 1;
}
// some more math code.

一些线程可能会根据条件进入其中一个，一些线程可能会进入两者，而其他线程可能不会进入其中任何一个。

现在，为了使所有线程在条件满足后再次回到“做同样的事情”，我应该在条件满足后使用__syncthreads()？或者这会以某种方式自动发生吗？
两个线程可以吗not由于其中一个操作落后而做同样的事情，从而毁了每个人？或者是否有一些幕后努力让他们在分支之后再次做同样的事情？

在扭曲内，任何线程都不会“领先于”任何其他线程。如果存在一个条件分支，并且它被 warp 中的某些线程采用，但其他线程没有采用（也称为 warp“发散”），则其他线程将闲置，直到分支完成，并且它们都在公共指令上“聚合”在一起。因此，如果您只需要线程的扭曲内同步，那么就会“自动”发生。

但不同的扭曲不会以这种方式同步。因此，如果您的算法要求某些操作在多个 warp 上完成，那么您将需要使用显式同步调用（请参阅 CUDA 编程指南，第 5.4 节）。

EDIT:重新组织了接下来的几段以澄清一些事情。

这里实际上有两个不同的问题：指令同步和内存可见性。

• __syncthreads()强制指令同步并确保内存可见性，但仅限于块内，而不是跨块（CUDA 编程指南，附录 B.6）。它对于共享内存上的先写后读很有用，但不适合同步全局内存访问。

• __threadfence()确保全局内存可见性，但不执行任何指令同步，因此根据我的经验，它的用途有限（但请参阅附录 B.5 中的示例代码）。

• 内核中不可能进行全局指令同步。如果你需要f()在调用之前在所有线程上完成g()在任何线程上，拆分f() and g()分成两个不同的内核并从主机串行调用它们。

• 如果您只需要增加共享或全局计数器，请考虑使用原子增量函数atomicInc()（附录 B.10）。对于上面的代码，如果x1 and x2不是全局唯一的（在网格中的所有线程中），非原子增量将导致竞争条件，类似于附录 B.2.4 的最后一段。

最后，请记住，对全局内存的任何操作，特别是同步函数（包括原子）都会损害性能。

在不知道您要解决的问题的情况下，很难推测，但也许您可以重新设计算法，在某些地方使用共享内存而不是全局内存。这将减少同步的需要并提高性能。