optimization - CUDA:同步线程

Question

几乎在我读到的有关使用 CUDA 编程的任何地方都提到了经线中的所有线程都做同样事情的重要性。
在我的代码中，我无法避免某种情况。它看起来像这样:

// 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() 同步它们吗？ ?或者这是否会以某种方式自动发生？
两个线程会不会因为其中一个线程落后于一个操作而不能做同样的事情，从而对每个人都毁了它？或者是否有一些幕后努力让他们在分支后再次做同样的事情？

Answer 1

在经纱内，没有任何线程会“领先”于其他线程。如果有一个条件分支，并且它被 warp 中的一些线程而不是其他线程占用(又名 warp“发散”)，则其他线程将一直空闲，直到分支完成，并且它们都“收敛”到一个共同的指令上.因此，如果您只需要线程的 Warp 内同步，那会“自动”发生。

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

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

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

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

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

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

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

最后，请记住，对全局内存的任何操作，尤其是同步函数(包括原子)对性能都是不利的。

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