同一块中的扭曲如何发散

Question

我有点困惑 Warps 如何可能发散并需要通过

__syncthreads()

函数进行同步。块中的所有元素以 SIMT 方式处理相同的代码。他们怎么可能不同步呢？和调度器有关系吗？不同的扭曲有不同的计算时间吗？为什么使用

__syncthreads()

时会有开销？

假设我们在一个块中有 12 个不同的经线，其中 3 个已经完成了工作。所以现在有空闲，其他扭曲得到它们的计算时间。或者他们仍然有计算时间来执行

__syncthreads()

功能吗？

Answer 1

首先让我们注意术语。 Warp 发散是指由于代码中的控制结构（if、while 等），单个 warp 内的线程采用不同的执行路径。您的问题确实与 warp 和 warp 调度有关。 尽管 SIMT 模型可能建议

所有

线程以锁步方式执行，但事实并非如此。 首先，不同块内的线程是完全独立的。 它们可以按彼此之间的任何顺序执行。 对于关于同一块内线程的问题，我们首先观察一个块最多可以有 1024 个线程，但是今天的 SM（SM 或 SMX 是处理线程块的 GPU 内部的“引擎”）没有 1024 个 cuda 核心，因此理论上 SM 不可能以锁步的方式执行线程块的所有线程。 请注意，单个线程块在单个 SM 上执行，而不是同时在所有（或多个）SM 上执行。 因此，即使一台机器总共有 512 个或更多 cuda 核心，它们也不能全部用于处理单个线程块的线程，因为单个线程块在单个 SM 上执行。 （这样做的原因之一是，线程块中的所有线程都可以访问 SM 特定的资源，例如共享内存。） 那么会发生什么呢？ 事实证明，每个 SM 都有一个 warp 调度程序。 warp只不过是 32 个线程的集合，这些线程被分组在一起、一起调度并一起执行。 如果一个线程块有 1024 个线程，那么它有 32 个线程束，每个线程束有 32 个线程。 例如，现在，在 Fermi 上，SM 有 32 个 CUDA 核心，因此可以合理地考虑 SM 以锁步方式执行扭曲（这就是 Fermi 上发生的情况）。 通过锁步，我的意思是（忽略经纱发散的情况，以及指令级并行性的某些方面，我试图在这里保持简单的解释......）在上一条指令之前，经纱中没有指令被执行已被 warp 中的所有线程执行。 因此，Fermi cc2.0 SM 实际上只能在任何给定时刻发出线程块中的一个扭曲。 该线程块中的所有其他扭曲都已排队，准备就绪，等待。

现在，当一个扭曲的执行因任何原因而停滞时，扭曲调度程序可以自由地将该扭曲移出并引入另一个准备好的扭曲（这个新的扭曲甚至可能不是来自同一个线程块，但是我离题了。）希望现在您可以看到，如果一个线程块中包含超过 32 个线程，则并非所有线程实际上都在同步执行。 一些经线在其他经线之前进行。 这种行为通常是可取的，除非并非如此。 有时，您不希望线程块中的any线程继续超出某个点，直到满足条件。 这就是

__syncthreads()

的用途。 例如，您可能正在将数据从全局复制到共享内存，并且您不希望在正确填充共享内存之前开始任何线程块数据处理。

__syncthreads()

确保所有线程都有机会复制其数据元素，然后任何线程可以超越屏障并可能开始对现在驻留在共享内存中的数据进行计算。

__syncthreads()

的开销有两种形式。 首先，处理与此内置函数相关的机器级指令的成本非常小。 其次，

__syncthreads()

通常会强制warp调度程序和SM对线程块中的所有warp进行洗牌，直到每个warp遇到障碍。 如果这有用，那就太好了。 但如果不需要，那么你就是在花时间做一些不需要的事情。 因此，建议不要只是在代码中随意撒上

__syncthreads()

。 在需要的地方谨慎使用它。 如果您可以设计一种不像其他算法那样频繁使用它的算法，那么该算法可能会更好（更快）。