我有 GeForce GTX460 SE，所以它是：6 SM x 48 CUDA Cores = 288 CUDA Cores。众所周知，一个 Warp 中包含 32 个线程，并且在一个块中同时（一次）只能执行一个 Warp。也就是说，在单个多处理器（SM）中只能同时执行一个 Block、一个 Warp 和只有 32 个线程，即使有 48 个内核可用？
此外，可以使用threadIdx.x和blockIdx.x来分配具体的Thread和Block的示例。要分配它们，请使用内核 <<< Blocks, Threads >>> ()。但是如何分配特定数量的 Warp 并分配它们，如果不可能，那么为什么还要费心去了解 Warp 呢？
情况比你描述的要复杂得多。
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 调度器都会尝试“配对”两个 warp 来调度，以最大化 SM 的利用率。
如果指令是独立的，warp 调度程序可以从不同的块或同一块中的不同位置发出 warp。因此，可以同时处理来自多个块的扭曲。
更复杂的是，执行指令的资源少于 32 个的 warp 必须多次发出，以便为所有线程提供服务。例如，有 8 个 SFU，这意味着包含需要 SFU 的指令的 warp 必须被调度 4 次。
这个描述被简化了。还有其他一些限制也在起作用，它们决定了 GPU 如何安排工作。您可以通过在网上搜索“fermi architecture”找到更多信息。
所以，来到你的实际问题，
为什么要费心去了解 Warps？
当您尝试最大化算法的性能时，了解 warp 中的线程数并将其考虑在内变得很重要。如果您不遵守这些规则，您将失去性能：
在内核调用中<<<Blocks, Threads>>>，尝试选择与 warp 中的线程数均分的线程数。如果你不这样做，你最终会启动一个包含非活动线程的块。
在您的内核中，尝试让 warp 中的每个线程都遵循相同的代码路径。如果你不这样做，你就会得到所谓的经线发散。发生这种情况是因为 GPU 必须通过每个不同的代码路径运行整个 warp。
在您的内核中，尝试让每个线程都以特定模式加载和存储数据。例如，让 warp 中的线程访问全局内存中的连续 32 位字。
线程是否必须按顺序分组为 Warps，1 - 32、33 - 64 ...？
是的，编程模型保证线程按特定顺序分组到 warp 中。
作为优化发散代码路径的一个简单示例，可以使用将块中的所有线程分成 32 个线程组吗？例如：switch (threadIdx.s/32) { case 0: /* 1 warp*/ break; case 1: /* 2 warp*/ break; /* ETC */ }
确切地 ：）
单个 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 中缓存数据或最终处理多少事务。