如何测量 CUDA 中的流式多处理器使用/空闲时间？

Question

一个简单的问题，真的：我有一个内核，它以每个流式多处理器 (SM) 可能的最大块数运行，我想知道理论上我可以从中提取多少性能。理想情况下，我想知道空闲的 SM 周期的百分比，即所有 warp 在内存访问时都被阻塞。

我真的只是想找到那个号码。我不在寻找的是

• 关于增加入住率的一般提示。我正在使用我能获得的所有占用率，即使我设法获得更高的性能，它也不会告诉我理论上还有多少可能。
• 如何计算理论峰值 GFlops。我的计算不是以 FP 为中心的，还有很多整数算术和逻辑也在进行。

Answer 1

Nsight Visual Studio 版本 2.1 和 2.2 问题效率实验提供您请求的信息。这些计数器/指标应在 CUDA 5.0 之后的版本中添加到 Visual Profiler。

Nsight Visual Studio 版

来自Nsight Visual Studio Edition 2.2 User Guide | Analysis Tools | Other Analysis Reports | Profiler CUDA Settings | Issue Efficiency Section

问题效率提供有关设备能力的信息 发出内核的指令。报告的数据包括 执行依赖关系、合格的 warp 和 SM 停止原因。

对于计算能力为 2.x 的设备，一个多处理器有两个 warp 调度程序。每个 warp 调度器最多管理 24 个 warp，总共 每个多处理器 48 个扭曲。内核执行配置可能 减少运行时间限制。有关入住的信息，请参阅 完成占用实验。第一个调度器负责 带有奇数 ID 的 warp，第二个调度器负责 warp 具有偶数 ID。

KEPLER 更新：对于计算能力 3.x，多处理器有四个 warp 调度程序。每个 warp 调度程序最多管理 16 个 warp，每个多处理器总共管理 64 个 warp。

在每个指令发出时间，每个调度程序都会选择一个符合条件的 warp 从它的活动 warp 列表中删除并发出指令。经线是 如果指令已被提取，则执行单元符合条件 指令要求的可用，而指令没有 尚未满足的依赖项。

调度器报告以下关于warp的统计信息 多处理器：

Active Warps – 一个 Warp 从它被安排在一个 多处理器，直到它完成最后一条指令。活跃的 扭曲计数器每个周期递增 0-48。每最大增量 周期由理论占用率定义。

KEPLER UPDATE 范围是每个周期 0-64。

Eligible Warps – 如果一个活动的 Warp 能够发出下一条指令，则它是合格的。 不符合条件的 Warp 将报告问题停顿原因。这 计数器将在每个周期递增 0-ActiveWarps。

更新在 Fermi 上，Issue Stall Reason 计数器仅在 warp 调度程序没有符合条件的 warp 的周期上更新。在 Kepler 上，即使 warp 调度程序发出指令，问题停顿原因计数器也会在每个周期更新。

Zero Eligible Warps – 如果出现以下情况，此计数器将每个周期递增 1 调度程序都没有 可以发出的warp。

一个合格的 Warp – 此计数器递增 如果两个调度程序中只有一个具有可以 发布。

KEPLER 更新：在 Kepler 上，计数器是每个调度程序的，因此 One Eligible Warp 意味着 warp 调度程序可以发出指令。在 Fermi 上，两个调度程序都有一个计数器，因此在 Fermi 上，您希望 One Eligible Warp 计数器尽可能小。

Warp Issue Holes – 此计数器在每个周期递增 不符合条件的活动经线数量。这与 Active Warps 减去 Eligible Warps。

长翘曲问题孔 – 这个 计数器在每个周期增加一个活跃的经纱的数量 没有资格发出指令超过 32 个时钟 循环。长孔表示扭曲在长延迟时停止 障碍和内存操作等原因。

问题停顿原因 - 每个周期每个不合格的扭曲都会增加一个问题停顿 原因计数器。所有问题停止原因计数器的总和相等 扭曲问题孔。不合格的扭曲将增加指令获取 如果下一条汇编指令尚未执行，则停止原因 获取。 Execution Dependency 如果输入依赖项是 尚不可用。这可以通过增加数量来减少 独立指令。 数据请求请求停止的原因 由于所需资源不可用，目前无法制作， 或已充分利用，或该类型的操作已经太多 杰出的。如果数据请求占很大一部分 失速原因，还应该运行内存实验来确定 如果您可以根据请求优化现有事务，或者如果您需要 重新审视你的算法。 纹理如果纹理停滞原因 子系统已被充分利用，目前无法接受 进一步的操作。 Synchronization 停止原因，如果 warp 是 在 __syncthreads() 处阻塞。如果这个原因很大而且内核 执行配置仅限于少量块然后 考虑将内核网格划分为更多的线程块。

视觉分析器 5.0

Visual Profiler 没有解决您问题的计数器。在添加计数器之前，您可以使用以下计数器：

• sm_efficiency[_instance]
• ipc[_instance]
• achieved_occupancy。

计算能力的目标和最大 IPC 是：

Compute     Target   Max
Capability  IPC      IPC
2.0         1.7      2.0
2.x         2.3      4.0
3.x         4.4      7.0

目标 IPC 用于 ALU 受限计算。内存绑定内核的目标 IPC 会更少。对于计算能力 2.1 及更高版本的设备，使用 IPC 更加困难，因为每个 warp 调度程序都可以同时发出。