【发布时间】：2014-07-10 01:40:28
【问题描述】：
我对在 GPU 上运行的内核的吞吐量有疑问。假设它的占用率为 0.5，块大小为 256：编程指南指出最好有很多块，这样它们可以隐藏内存延迟等。但我不明白为什么这是正确的。因为一旦内核每个 Streaming Multi-processor = 24 的 warp 数，即 3 个块，它将达到峰值吞吐量。因此，拥有超过 24 个 warp（或 3 个块）不会对吞吐量产生任何影响。
我错过了什么吗？谁能纠正我？
【问题讨论】：
【发布时间】：2014-07-10 01:40:28
【问题描述】：
虽然低占用率的 SM 确实无法充分隐藏延迟，但了解这一点很重要：
更高的入住率！= 更高的吞吐量！
占用率只是衡量 SM 在任何给定时刻可供选择的工作量。拥有更多常驻 warp 使 SM 更有能力做有用的工作，而其他 warp 正在等待结果（内存访问或计算的结果 - 两者都具有非零延迟）。
吞吐量是衡量每秒完成多少工作的指标，虽然它可能受到延迟（因此占用率）的限制，但它也可能受到内存带宽、指令吞吐量（执行单元的数量）和其他因素。
编程指南指出拥有多个线程块比只有一个大线程块更好的原因是因为有时不仅能够从其他线程块发出工作，而且还能够从其他块发出工作。这是一个例子：
假设您的大线程块必须从全局内存（高延迟）加载数据并将其存储到共享内存（低延迟），然后必须立即执行__syncthreads()。在这种情况下，当一个 warp 完成加载其数据并将其写入共享内存时，它必须等待直到块中的所有其他线程完成相同的操作。对于一个大街区来说，这可能需要相当长的时间。但是如果有多个较小的线程块占用 SM，那么 SM 可以在等待第一个块中满足__syncthreads 的同时从其他块切换并开始工作。这有助于减少 GPU 空闲时间并提高效率。
您不一定希望拥有非常小的块（因为 Fermi 上的 SM 最多支持 8 个常驻块），但拥有 128-512 个线程的块通常比使用具有 1024 个线程的块更有效。
【讨论】：
-
好的，你说的我明白了。这一切都是正确的，我相信。但是缺少一些东西。请使用我的示例：occupancy=0.5，blocksize=256。我觉得这对于许多内核来说很常见。这意味着我将在 GTX580 中拥有 24 个活动扭曲（或 3 个活动块）。那么启动 3 个区块和 6 个区块的吞吐量有什么区别呢？造成差异的原因是什么？
-
更高的占用率！= 更高的吞吐量！因此，在不了解内核瓶颈的具体情况下，通常无法预测启动 3 块和 6 块之间的吞吐量差异。现在，如果占用率限制为 0.5（通过寄存器或共享内存使用），并且您的 GPU 至少有 6 个 SM，那么启动 6 个块而不是 3 个应该会使您的吞吐量翻倍，因为这会使用更多短信。通常，如果每个 SM 不是多个，您至少希望启动与 SM 一样多的块。但我不认为这是你要问的......
-
哦，对不起，我搞砸了。我想说的是每个 SM 3 个块，每个 SM 6 个块（不是块的总数）。是的，我很确定这将取决于内核的瓶颈，但是如何呢？而且应该有一些一般情况，对吧？ （假设占用率保持在 0.5 - 实际上它不会受到块数的影响）
-
或者说简单点，就说内核占用率为0.5。这意味着每个 SM 可以有 3 个活动块。在 GPU 中，有 16 个 SM。当我总共启动 48 个块时，IPC（每周期指令）是 100。但是当我总共启动 64 个块时，IPC 是 120，可能是什么原因造成这种差异？
-
我认为你的意思是 1.0 和 1.2 而不是 100 和 120，但我的猜测是你的块的启动时间有一些偏差，或者你的块有可变的运行时间，所以有更多在最初的 3 个区块中的一个或多个完成后，超过 48 个区块有助于填补 SM 上的一些空白，从而提高整体效率。
如果您的启用 cuda 的卡中只有 一个 SM，那么在您的情况下，拥有 3 个以上的块不会改变吞吐量。通常一个 GPU 中有 8 个或更多 SM。
此外，在一个 SM 上运行的块数不仅仅取决于 warp 的数量。这只是一个限制因素，还有许多其他因素。 CUDA Occupancy Calculator 是查看内核占用率的绝佳工具。
【讨论】：
-
那么，如果我的卡中有多个 SM，吞吐量会发生什么变化？说 16。我知道 Cuda 占用计算器，但在这里没有帮助。我要问的是“拥有足够的经线（你所拥有的只是活跃的经线）和拥有足够多的经线（你有一些活跃的经线，而其他经线是空闲的）有什么区别。我敢打赌答案不会这么短。