cuda - block ，线程，warpSize

Question

关于如何选择#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。 (不是一些 HWcore 运行 kernelA，而其他一些运行 kernelB。)因此，如果我只有一个线程要运行，我就是在“浪费”其他 HWcore。真的？

3)如果warp-scheduler以warpSize为单位发出工作(32 个线程)，并且每个 SM 有 32 个 HWcores，那么 SM 将被充分利用。当 SM 有 48 个 HWcore 时会发生什么？当调度程序以 32 个块为单位发出工作时，如何让所有 48 个内核都得到充分利用？ (如果上一段是真的，那么调度程序以 HWcore 大小为单位发布工作不是更好吗？)

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

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

5a) 我认为越小越好，但我做的多小有关系nBlocks ?问题是，如何选择nBlocks的值那是体面的吗？ (不一定是最优的。)是否有选择 nBlocks 的数学方法？ ，或者只是试错。

Answer 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 的扭曲仅增加了 33%。 Nsight VSE 问题效率实验是确定是否有足够的扭曲来隐藏指令和内存延迟的最佳方法。不幸的是，Visual Profiler 没有这个指标。

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

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