13

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

4

4 回答 4

36

在一个经线中，没有任何线程会“领先于”任何其他线程。如果有一个条件分支并且它被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 的最后一段。

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

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

于 2009-10-29T17:07:03.860 回答
2

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

任何流控制指令（if、switch、do、for、while）都可以通过导致同一 warp 的线程发散来显着影响指令吞吐量；也就是说，遵循不同的执行路径。如果发生这种情况，则必须对不同的执行路径进行序列化，从而增加为此 warp 执行的指令总数。当所有不同的执行路径都完成后，线程会汇聚回相同的执行路径。

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

于 2009-10-29T16:57:15.943 回答
2

在加布里埃尔的回应中：

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

如果您在同一个线程中需要 f() 和 g() 的原因是因为您正在使用寄存器内存，并且您希望从 f 到 g 的寄存器或共享数据？也就是说，对于我的问题，跨块同步的全部原因是因为 g 中需要来自 f 的数据 - 并且突破到内核将需要大量额外的全局内存来将寄存器数据从 f 传输到 g，我想避免

于 2012-12-07T22:34:14.617 回答
1

你的问题的答案是否定的。你不需要做任何特别的事情。无论如何，您可以解决此问题，而不是您的代码，您可以执行以下操作：

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

您应该检查是否可以使用共享内存并以合并模式访问全局内存。还要确保您不想在超过 1 个线程中写入同一个索引。

于 2010-02-09T00:05:50.083 回答