为啥 CUDA GPU 只需要 8 个活动 warp?

Posted

技术标签:

【中文标题】为啥 CUDA GPU 只需要 8 个活动 warp?【英文标题】:Why does CUDA GPU only need 8 active warps?为什么 CUDA GPU 只需要 8 个活动 warp? 【发布时间】:2015-05-29 09:31:52 【问题描述】:

如this work中所说:

如果 CUDA 编译器生成的指令流表示 ILP 为 3.0(即平均可以执行 3 条指令之前发生危险),并且指令流水线深度为 22 阶段,则只有 8 个活动 warp( 22 / 3) 可能足以完全隐藏指令延迟并实现最大算术吞吐量。

我不明白为什么就足够了?

【问题讨论】:

【参考方案1】:

如果调度程序可以在每个指令发出周期成功地从同一个 warp 发出一条指令,持续 22 个连续的周期,那么调度程序没有理由调度另一个 warp 来代替它,并且单个 warp 足以填充管道。这将对应于至少 22 的 ILP。

但 Real-World Code™ 从未表现出如此高的 ILP:例如,某些指令取决于先前指令的结果或内存请求。当调度器不能再执​​行独立指令时,warp 的执行就会停止。调度程序将选择另一个准备执行的 warp,并执行尽可能多的指令,直到该 warp 也停止,依此类推。

因此,如果 warp #1 成功执行 3 条指令然后停止,调度器选择 warp #2,执行 3 条指令......等等。当调度器到达 warp #8 时,管道中已经有 21 条指令用于 7停滞的扭曲。然后从该扭曲执行一条指令就足以完全填满管道。当管道开始耗尽时,warp #1 再次准备就绪,因此为什么 ILP 为 3 的 8 个 warp 足以填充 22 阶段的管道。

【讨论】:

以上是关于为啥 CUDA GPU 只需要 8 个活动 warp?的主要内容,如果未能解决你的问题,请参考以下文章

CUDA:为啥会有大量的 GPU 空闲时间?

为啥启动 Numba cuda 内核最多可使用 640 个线程,但在有大量可用 GPU 内存时却因 641 而失败?

为啥允许我运行块数超过 GPU 的 CUDA 核心数的 CUDA 内核?

用于在 GPU 上监视 CUDA 活动的类***实用程序

为啥同时使用 numba.cuda 和 CuPy 从 GPU 传输数据这么慢?

为啥在使用 cudaMallocManaged 时 NVIDIA Pascal GPU 运行 CUDA 内核速度很慢