块，线程，warpSize

Posted 2023-04-15

【中文标题】块，线程，warpSize

【英文标题】：blocks, threads, warpSize

【发布时间】：2012-06-11 15:37:12

【问题描述】：

关于如何选择#blocks 和blockSize 的讨论很多，但我仍然遗漏了一些东西。我的许多担忧都解决了这个问题：How CUDA Blocks/Warps/Threads map onto CUDA Cores?（为了简化讨论，有足够的 perThread 和 perBlock 内存。内存限制在这里不是问题。）

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal);

1) 为了让 SM 尽可能忙碌，我应该将 nThreads 设置为 warpSize 的倍数。是吗？

2) 一个 SM 一次只能执行一个内核。也就是说，该 SM 的所有 HWcore 都只执行 kernelA。 （不是一些运行 kernelA 的 HWcore，而其他运行 kernelB。）因此，如果我只有一个线程要运行，我将“浪费”其他 HWcore。是吗？

3) 如果 warp-scheduler 以warpSize（32 个线程）为单位发出工作，并且每个 SM 有 32 个 HWcore，则 SM 将被充分利用。当 SM 有 48 个 HWcore 时会发生什么？当调度程序以 32 个块的形式发布工作时，如何保持所有 48 个内核的充分利用？ （如果上一段是真的，调度器以HWcore大小为单位下发工作不是更好吗？）

4) 看起来 warp-scheduler 一次将 2 个任务排队。这样当当前正在执行的内核停顿或阻塞时，就会换入第二个内核。（不清楚，但我猜这里的队列深度超过 2 个内核。）这是正确的吗？

5) 如果我的硬件的上限为每块 512 个线程 (nThreadsMax)，这并不意味着拥有 512 个线程的内核将在一个块上运行得最快。 （同样，内存不是问题。）如果我将 512 线程内核分布在许多块上，而不仅仅是一个块，我很有可能会获得更好的性能。该块在一个或多个 SM 上执行。是吗？

5a) 我认为越小越好，但我把nBlocks 做得有多小有关系吗？问题是，如何选择nBlocks的值才算体面？ （不一定是最优的。）是否有选择nBlocks 的数学方法，或者只是简单地试错。

【问题讨论】：

这个 GPU 有 192 个 CudaCore。那将是 4 个 SM，具有 48 个硬件核心 (HWcores)。 4 * 48 = 192

【参考方案1】：

1) 是的。

2) CC 2.0 - 3.0 设备最多可以同时执行 16 个网格。每个 SM 限制为 8 个块，因此为了达到完全并发，设备必须至少有 2 个 SM。

3) 是的，warp 调度程序会在时间选择和发出 warp。忘记 CUDA 内核的概念，它们无关紧要。为了隐藏延迟，您需要具有高指令级并行性或高占用率。对于 CC 1.x，建议 >25%，对于 CC >= 2.0，建议 >50%。一般来说，CC 3.0 需要比 2.0 设备更高的占用率，因为调度器增加了一倍，但每个 SM 的 warp 只增加了 33%。 Nsight VSE 问题效率实验是确定您是否有足够的扭曲来隐藏指令和内存延迟的最佳方法。不幸的是，Visual Profiler 没有这个指标。

4) warp 调度器算法没有文档化；但是，它不考虑线程块源自哪个网格。对于 CC 2.x 和 3.0 设备，CUDA 工作分配器将在分配来自下一个网格的块之前分配来自一个网格的所有块；但是，编程模型无法保证这一点。

5) 为了让 SM 保持忙碌，您必须有足够的块来填充设备。之后，您要确保有足够的经线来达到合理的占用率。使用大线程块有利有弊。大线程块通常使用较少的指令缓存并且在缓存上的占用空间较小；然而，大线程块在同步线程处停滞（SM 可能变得不那么有效，因为可供选择的扭曲更少）并且倾向于让指令在相似的执行单元上执行。我建议每个线程块尝试 128 或 256 个线程来启动。较大和较小的线程块都有充分的理由。 5a) 使用占用计算器。选择太大的线程块大小通常会导致您受到寄存器的限制。选择太小的线程块大小会发现您受到共享内存或每个 SM 8 个块的限制。

【参考方案2】：

让我试着一一回答你的问题。

没错。 “HWcores”到底是什么意思？你陈述的第一部分是正确的。

根据NVIDIA Fermi Compute Architecture Whitepaper：“SM 以 32 个并行线程为一组调度线程，称为 warp。每个 SM 具有两个 warp 调度器和两个指令调度单元，允许同时发出和执行两个 warp。Fermi 的双 warp 调度器选择两个 warp，并从每个 warp 发出一条指令到一组 16 个核心、16 个加载/存储单元或四个 SFU。因为 warp 独立执行，Fermi 的调度程序不需要从指令流中检查依赖关系。

此外，NVIDIA Keppler Architecture Whitepaper 声明：“Kepler 的四次 warp 调度程序选择四个 warp，每个 warp 可以调度两个独立指令。”

因此，“多余的”核心用于一次调度多个 warp。

warp 调度程序调度同一个内核的warp，而不是不同的内核。

不完全正确：每个块都锁定在一个 SM 中，因为这是其共享内存所在的位置。 这是一个棘手的问题，取决于您的内核是如何实现的。您可能想看看 Vasily Volkov 的 nVidia 网络研讨会Better Performance at Lower Occupancy，它解释了一些更重要的问题。不过，首先，我建议您选择线程数以提高占用率，使用CUDA Occupancy Calculator。

【讨论】：

感谢您的回答。我看过占用计算器。它很有用，但也需要更新。它允许选择计算版本 2.0，但不允许 ThreadsPerBlock 超过 512（计算 1.x）。