cuda - 为什么要费心去了解 CUDA Warps？

Question

我有 GeForce GTX460 SE，所以它是:6 SM x 48 CUDA 核心 = 288 CUDA 核心。众所周知，一个Warp中包含32个线程，并且一个 block 中同时(一次)只能执行一个Warp。也就是说，在单个多处理器(SM)中，即使有 48 个可用核心，也只能同时执行 1 个 Block、1 个 Warp 和 32 个线程？

另外，分配具体Thread和Block的例子可以使用threadIdx.x和blockIdx.x。要分配它们，请使用 kernel <<< Blocks, Threads >>> ()。但是如何分配特定数量的 Warp 并分发它们，如果不可能那么为什么还要费心去了解 Warp呢？

Answer 1

情况比你描述的要复杂得多。

ALU(核心)、加载/存储 (LD/ST) 单元和特殊功能单元 (SFU)(图中绿色)是流水线单元。它们在完成的各个阶段同时保留许多计算或操作的结果。因此，在一个周期中，它们可以接受一项新操作，并提供很久以前开始的另一项操作的结果(如果我没记错的话，ALU 大约有 20 个周期)。因此，理论上单个 SM 拥有同时处理 48 * 20 个周期 = 960 个 ALU 操作的资源，即每个 warp 960/32 个线程 = 30 个 warp。此外，它还可以处理 LD/ST 操作和 SFU 操作，无论其延迟和吞吐量如何。

warp 调度程序(图中黄色)可以在每个周期为每个 warp 调度 2 * 32 个线程 = 64 个线程到管道。这就是每个时钟可以获得的结果数。因此，考虑到存在混合的计算资源，48 个核心、16 个 LD/ST、8 个 SFU，每个资源都有不同的延迟，因此同时处理混合的扭曲。在任何给定的周期，warp 调度程序都会尝试“配对”两个 warp 进行调度，以最大限度地提高 SM 的利用率。

如果指令是独立的，warp 调度程序可以从不同的 block 或同一 block 中的不同位置发出 warp。因此，可以同时处理来自多个 block 的扭曲。

执行资源少于 32 个指令的 warp 必须多次发出才能服务所有线程，这进一步增加了复杂性。例如，有 8 个 SFU，这意味着包含需要 SFU 的指令的 warp 必须被调度 4 次。

此描述已简化。还有其他限制也决定了 GPU 如何调度工作。您可以通过在网络上搜索“费米架构”来找到更多信息。

那么，回到你的实际问题，

why bother to know about Warps?

当您尝试最大限度地提高算法性能时，了解扭曲中的线程数量并将其考虑在内变得很重要。如果您不遵守这些规则，您的性能就会下降:

• 在内核调用中，<<<Blocks, Threads>>> ，尝试选择与经纱中的线程数均匀划分的线程数。如果不这样做，您最终会启动一个包含非事件线程的 block 。

• 在内核中，尝试让 warp 中的每个线程遵循相同的代码路径。如果不这样做，就会得到所谓的扭曲发散。发生这种情况是因为 GPU 必须通过每个不同的代码路径运行整个扭曲。

• 在内核中，尝试让每个线程进行扭曲加载并以特定模式存储数据。例如，让 warp 中的线程访问全局内存中的连续 32 位字。