CUDA：同步线程

几乎在任何地方，我读过关于使用CUDA编程的内容，都提到重要性，即变形中的所有线程都执行相同的操作。
在我的代码中，我有一种情况，我无法避免某种情况。它看起来像这样：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()之后同步它们吗？或者这是以某种方式自动发生的？
可以两个线程是而不是做同样的事情，由于其中一个是后面的一个操作，从而毁灭了每个人？还是有一些幕后的努力让他们在分支之后再次做同样的事情？

来源

2009-10-29 shoosh

在变形中，没有线程会“超前”任何其他线程。如果有一个条件分支，并且它是由warp中的某些线程获取的，但不是其他线程（又称warp“divergence”），则其他线程将空闲直到分支完成，并且它们全部“聚合”回到一个共同指令。所以如果你只需要线程内部的同步，那就会“自动地”发生。

但是不同的经纱不是以这种方式同步的。因此，如果您的算法要求某些操作在许多经线上完成，那么您需要使用显式同步调用（请参见CUDA编程指南，第5.4节）。

编辑：重组未来数段，以澄清一些事情。

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

__syncthreads()强制指令同步，确保存储器的可见性，但只能在一个块，而不是跨越块（CUDA编程指南附录B.6）。这对于在共享内存上执行写入然后读取很有用，但不适用于同步全局内存访问。
__threadfence()确保全局内存可见性但不做任何指令同步，因此根据我的经验，它的使用有限（但请参阅附录B.5中的示例代码）。
全局指令同步是不可能的内核中。如果您在任何线程调用g()之前需要f()上的所有线程完成，拆分f()和g()成两个不同的内核，并从主机顺序调用它们。
如果您只需增加共享或全局计数器，请考虑使用原子增量函数atomicInc()（附录B.10）。对于上面的代码，如果x1和x2不是全局唯一的（跨网格中的所有线程），非原子增量将导致竞争条件，类似于附录B.2.4的最后一段。

最后，请记住，对全局存储器和特别是同步功能（包括原子）的任何操作都不利于性能。

不知道你解决它的问题是很难猜测，但也许你可以重新设计你的算法使用共享内存而非全局内存在一些地方。这将减少同步的需要并为您提供性能提升。

来源

2009-10-29 17:07:03 Gabriel

让我看看我是否明白。所以如果条件分支做相同数量的工作，这应该不会影响性能，因为每个线程都不会空闲太久。我对吗？ – 2011-03-31 22:51:03

@omegatai我认识到你的评论是旧的，但其他人可能想知道，所以在这里：一个warp一次只能处理一条指令，所以如果warp中的某些线程做了一件事，其余的线程做其他事情，总时间是这两组线程的总时间量。在经纱内没有时间重叠。性能受到影响。 – 2013-10-15 21:11:31

从CUDA最佳实践指南的第6.1节：

任何流控制指令（如果，开关，做，对，而）可以显著通过使相同的经纱的纱线以影响的指令吞吐量发散;即遵循不同的执行路径。如果发生这种情况，必须对不同的执行路径进行串行化处理，从而增加为此 warp执行的指令总数。当所有不同的执行路径都完成后，线程会聚到回到相同的执行路径。

所以，你不需要做任何特别的事情。

来源

2009-10-29 16:57:15 Edric

您的问题的答案是否定的。你不需要做任何特别的事情。无论如何，你可以解决这个问题，而不是你的代码，你可以做这样的事情：

buffer[x1] += (d1 < 0.5); 
buffer[x2] += (d2 < 0.5);

你应该检查是否可以在聚结的图案使用共享内存和访问全局内存。另外请确保您不想在多个线程中写入相同的索引。

来源

2010-02-09 00:05:50 purpletentacle

诀窍很微妙，但通过这样做，你根本不会分支！ – purpletentacle 2014-11-14 16:13:04

在加布里埃尔的回应：

“环球指令同步是不可能的内核中如果需要调用任何线程克（），分割F（）和g之前所有线程完成F（）（）分成两个不同的内核，并从主机串行调用它们。“

如果你需要什么？在同一个线程F（）和g（）的原因是因为你使用的寄存器内存，你想从F寄存器或共享数据让到g？也就是说，对于我的问题，跨块同步的全部原因是因为f中的数据需要在g中 - 而向内核发展则需要大量额外的全局内存将寄存器数据从f传输到g，我想避免

来源

2012-12-07 22:34:14

CUDA：同步线程

回答

相关问题