5

关于如何选择#blocks & blockSize 已经有很多讨论，但我仍然遗漏了一些东西。我的许多担忧都解决了这个问题：CUDA Blocks/Warps/Threads 如何映射到 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，并且每个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，或者只是试错法。

4

2 回答 2

5

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) 使用占用计算器。选择太大的线程块大小通常会导致您受到寄存器的限制。

于 2012-06-07T20:10:17.420 回答
3

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

  1. 那是对的。
  2. “HWcores”到底是什么意思？你陈述的第一部分是正确的。

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

    此外，NVIDIA Keppler 架构白皮书指出：“Kepler 的四次 warp 调度程序选择四个 warp，每个 warp 可以在每个周期调度两个独立的指令。”

    因此，通过一次调度一个以上的 warp 来使用“多余的”核心。

  4. warp 调度器调度相同内核的warp，而不是不同的内核。

  5. 不完全正确：每个块都被锁定在一个 SM 中，因为那是它的共享内存所在的位置。
  6. 这是一个棘手的问题，取决于您的内核是如何实现的。您可能想看看 Vasily Volkov 的 nVidia Webinar Better Performance at Lower Occupancy，它解释了一些更重要的问题。不过，首先，我建议您使用CUDA Occupancy Calculator选择线程数以提高占用率。
于 2012-06-07T15:29:58.780 回答