为啥要费心去了解 CUDA Warps？

Posted 2023-03-23

【中文标题】为啥要费心去了解 CUDA Warps？

【英文标题】：Why bother to know about CUDA Warps?为什么要费心去了解 CUDA Warps？

【发布时间】：2012-08-02 17:53:36

【问题描述】：

我有 GeForce GTX460 SE，所以它是：6 SM x 48 CUDA 核心 = 288 CUDA 核心。 众所周知，一个 Warp 包含 32 个线程，而在一个块中同时（一次）只能执行一个 Warp。 即在单个多处理器（SM）中只能同时执行一个Block、一个Warp和只有32个线程，即使有48个内核可用？

此外，可以使用threadIdx.x和blockIdx.x来分配具体的Thread和Block的示例。要分配它们，请使用内核 >> ()。 但是如何分配特定数量的 Warp-s 并分发它们，如果不可能，那为什么还要费心去了解 Warp？

【问题讨论】：

您问题的第一段大部分内容完全不正确，因此您的问题的其余部分没有多大意义。

【参考方案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 调度程序可以从不同的块或同一块中的不同位置发出 warp。因此，可以同时处理来自多个块的扭曲。

更复杂的是，执行指令的资源少于 32 个的 warp 必须多次发出以供所有线程服务。例如，有 8 个 SFU，这意味着包含需要 SFU 的指令的 warp 必须被调度 4 次。

此描述已简化。还有其他一些限制也在起作用，它们决定了 GPU 如何安排工作。您可以通过在网上搜索“fermi architecture”找到更多信息。

所以，来到你的实际问题，

为什么要费心去了解 Warps？

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

在内核调用<<<Blocks, Threads>>> 中，尝试选择与warp 中的线程数均分的线程数。如果不这样做，您最终会启动一个包含非活动线程的块。

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

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

【讨论】：

谢谢，很好的回答！还有几个问题。 1.线程是否必须按顺序分组为 Warps，1 - 32、33 - 64 ...？ 2. 作为优化发散代码路径的一个简单示例，可以使用将块中的所有线程分隔为 32 个线程组吗？例如：switch (threadIdx.s/32) case 0: /* 1 warp*/ break; case 1: /* 2 warp*/ break; /* Etc */ 3. 单个 Warp 一次必须读取多少字节：4 字节 * 32 线程、8 字节 * 32 线程或 16 字节 * 32 线程？据我所知，一次对全局内存的一个事务接收 128 个字节。

【参考方案2】：

线程是否必须按顺序分组为 Warp，1 - 32、33 - 64 ...？

是的，编程模型保证线程按特定顺序分组到 warp 中。

作为优化发散代码路径的一个简单示例，可以使用将块中的所有线程分成 32 个线程组吗？例如：switch (threadIdx.s/32) case 0: /* 1 warp*/ break; case 1: /* 2 warp*/ break; /* 等等 */

没错:)

单个 Warp 一次必须读取多少字节：4 字节 * 32 线程、8 字节 * 32 线程或 16 字节 * 32 线程？据我所知，一次对全局内存的一个事务接收 128 个字节。

是的，全局内存的事务是 128 字节。因此，如果每个线程从连续地址读取一个 32 位字（它们可能也需要 128 字节对齐），那么 warp 中的所有线程都可以通过单个事务进行服务（4 字节 * 32 线程 = 128 字节）。如果每个线程读取的字节数更多，或者地址不连续，则需要发出更多事务（对每个单独的 128 字节行进行单独的事务处理）。

这在 CUDA 编程手册 4.2 的 F.4.2 节“全局内存”中有描述。那里还有一个宣传语，说情况与仅缓存在 L2 中的数据不同，因为 L2 缓存具有 32 字节的缓存行。我不知道如何安排数据仅在 L2 中缓存或最终处理多少事务。

【讨论】：

感谢您的澄清。对于仅在 L2 中缓存的数据，需要对 nvcc 使用编译器选项 -Xptxas -dlcm=cg。但我不知道在 VS 2010 中我必须在哪里写 (-Xptxas -dlcm=cg) :)

如果你能说出原子操作和 Warps。单个 Warp 的线程之间或一个块中不同 Warp 的线程之间的原子竞争（并发）哪个更好？我认为，当您访问共享内存时，一个经线的线程相互竞争比不同经线的线程少。并且相反地访问全局内存，一个块的不同warp的线程竞争比单个warp的线程少，不是吗？