我正在阅读 Cheng 等人的《专业 CUDA C 编程》。并且有一些示例说明如何运行（非常简单的单行）内核，例如，<<<1024, 512>>>性能比<<<2048, 256>>>. 然后他们（多次）声明您可能已经预料到了这个结果，因为第二次运行有更多的块，因此暴露了更多的并行性。但我不明白为什么。并行度不是由 SM 中并发扭曲的数量决定的吗？块大小与此有什么关系 - 这些扭曲属于哪个块并不重要 - 相同的块或不同的块，那么为什么使用较小的块会暴露更多的并行性（相反，如果块大小太小我会达到每个 SM 的最大块限制，从而导致并发扭曲更少）？我能想象的唯一场景是 Fermi 上的 1024 个线程块 = 32 个扭曲，每个 SM 限制最多有 48 个并发扭曲。这意味着只有 1 个并发块，并且只有 32 个并发扭曲是可能的，从而减少了并行量，但这是一个非常具体的用例。
更新：发布后我想到的另一件事是：在块中的所有扭曲完成之前，不能将其从 SM 中驱逐。因此，在该块执行结束时，可能会出现这样的情况：最后几个“最慢”的扭曲将整个块保留在 SM 中，而该块中的大部分扭曲已完成并停止，但不能再创建新的块。加载直到那几个正在执行的扭曲完成。所以这种情况下效率就变低了。现在，如果块更小，那么这种情况仍然会发生，但是相对于执行扭曲的停滞数量更小，因此效率更高。是这个吗？
是的，就是这样。你问题的第二段是一个很好的答案。
更详细地说，每个 SM 内的 warp 调度程序的数量是有限的（通常是 2 个）。每个 warp 调度程序都会跟踪许多活动的 warp，并且仅当允许 warp 在程序中进一步移动时才调度该 warp 的执行。Warp 调度程序跟踪的活动 Warp 数量有一个最大值（通常为 32）。因为在所有 warp 完成之前，线程块拥有的资源（例如共享内存）无法释放给新线程块，所以如果少数 warp 存在，那么大的块大小可能会导致调度程序可用的候选活动 warp 数量减少。需要很长时间才能完成。由于资源空闲或 SM 无法覆盖内存访问的延迟，这可能会导致性能下降。当使用或其变体之一跨线程块进行同步时，较大的块大小也会增加扭曲阻塞的可能性__syncthreads()，因此，可能会导致类似的现象。