cuda - NVIDIA GPU上的cuda Kernel的峰值吞吐量

Question

我对在GPU上运行的内核的吞吐量有疑问。假设它的占用率为0.5，块大小为256:编程指南指出，最好有许多块，以便它们可以隐藏内存延迟，等等。但是我不明白为什么这是正确的。因为一旦每个流多处理器内核具有24个扭曲数，即3个块，它将立即达到峰值吞吐量。因此，具有超过24个扭曲(或3个块)不会改变吞吐量。

我有什么想念的吗？谁能纠正我？

Answer 1

虽然低占用率SM确实不能充分掩盖延迟，但了解这一点很重要:

更高的占用率!=更高的吞吐量!

占用率只是对SM在任何给定时刻可供选择的工作量的一种度量。具有更多的常驻扭曲，使SM具有更多的能力来执行有用的工作，而其他扭曲却正在等待结果(内存访问或计算的结果-两者均具有非零延迟)。

吞吐量是衡量每秒完成多少工作的指标，尽管它可能会受到延迟(因此占用)的限制，但也可能会受到内存带宽，指令吞吐量(执行单元的数量)和其他因素的限制。

编程指南指出拥有多个线程块比仅包含一个大线程块更好的原因是，有时最好不仅能够从其他线程而且还可以从其他线程发出工作。这是一个例子:

想象一下，您的大线程块必须从全局内存中加载数据(高延迟)并将其存储到共享内存中(低延迟)，然后必须立即执行__syncthreads()。在这种情况下，当warp完成加载其数据并将其写入共享内存后，它必须等待直到块中的所有其他线程完成相同的操作。对于大块，可能要花相当长的时间。但是，如果有多个较小的线程块占用SM，则SM可以在等待第一个块中满足__syncthreads的同时切换并从其他块开始工作。这可以帮助减少GPU空闲时间并提高效率。

您不一定要有很小的块(因为Fermi上的SM支持最多8个驻留块)，但是拥有128-512个线程的块通常比使用1024个线程的块更有效。