CUDA:同步线程

Question

几乎在我读到有关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()吗？或者这是以某种方式自动发生的？
两个线程可以不做同样的事情,因为其中一个是后面的一个操作,因此毁掉了每个人吗？或者是否有一些幕后的努力让他们在分支后再次做同样的事情？

Answer 1

在经线内,任何线程都不会"超前".如果存在一个条件分支并且它由warp中的某些线程而不是其他线程(也就是warp"divergence")占用,则其他线程将空闲直到分支完成并且它们全部"聚合"在一起共同的指令.因此,如果您只需要线程内部同步,则会"自动"发生.

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

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

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

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

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

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

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

最后,请记住,对全局内存的任何操作,特别是同步函数(包括原子)都不利于性能.

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