我有点困惑 Warps 是如何发散并需要通过__syncthreads()函数同步的。Block 中的所有元素都以 SIMT 方式处理相同的代码。他们怎么可能不同步？它与调度程序有关吗？不同的经线是否有不同的计算时间？为什么使用时会有开销__syncthreads()？
假设我们在一个区块中有 12 个不同的 Warp，其中 3 个已经完成了他们的工作。所以现在有空闲，其他扭曲得到了他们的计算时间。还是他们仍然有计算时间来执行该__syncthreads()功能？
问问题
1424 次
1 回答
6
首先让我们小心术语。由于代码中的控制结构（if、while 等），warp分歧是指单个 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 以同步方式执行扭曲是合理的（即在费米上会发生什么）。通过锁步，我的意思是（忽略经线发散的情况，以及指令级并行性的某些方面，我试图在这里保持简单的解释......）在前一条指令之前没有执行经线中的指令已被warp中的所有线程执行。因此，Fermi SM 实际上只能在任何给定时刻执行线程块中的一个扭曲。该线程块中的所有其他经线都已排队，准备就绪，等待。
现在，当一个 warp 的执行由于任何原因停止时，warp 调度程序可以自由地将该 warp 移出并带入另一个准备就绪的 warp（这个新的 warp 甚至可能不是来自同一个线程块，但我离题了。）希望现在你可以看到如果一个线程块中有超过 32 个线程，并不是所有的线程实际上都在锁步中执行。一些经线在其他经线之前进行。
这种行为通常是可取的，除非它不是。有时您不希望线程块中的任何线程超过某个点，直到满足某个条件。这是__syncthreads()为了什么。例如，您可能正在将数据从全局复制到共享内存，并且您不希望在正确填充共享内存之前开始任何线程块数据处理。 __syncthreads()确保所有线程都有机会复制它们的数据元素，然后任何线程可以继续越过障碍，并且可能开始对现在驻留在共享内存中的数据进行计算。
开销有__syncthreads()两种口味。首先，处理与此内置函数相关的机器级指令的成本非常低。其次，__syncthreads()通常会强制 warp 调度程序和 SM 对线程块中的所有 warp 进行洗牌，直到每个 warp 遇到障碍。如果这有用，那就太好了。但是，如果不需要，那么您就是在花时间做一些不需要的事情。因此，建议不要随意散布__syncthreads()您的代码。在需要的地方谨慎使用它。如果您可以设计出一种不像其他算法那样使用它的算法，那么该算法可能会更好（更快）。
于 2012-11-23T17:00:53.597 回答