性能分析指标:吞吐量、延迟、占用率与内存带宽利用率

在进入 Nsight 工具实战之前,必须先建立一套统一的“度量语言”。GPU 性能分析不是简单地看某个计数器高或低,而是通过一组相互关联的指标,定位瓶颈究竟在计算单元、访存路径、调度机制还是并行度不足。本章围绕四个核心指标展开:吞吐量(Throughput)、延迟(Latency)、占用率(Occupancy)与内存带宽利用率(Memory Bandwidth Utilization)。理解它们的定义、测量方式与相互制约关系,是后续使用 Nsight Systems 和 Nsight Compute 的基础。

3.1 核心原理:四个指标如何描述 GPU 行为

GPU 设计哲学是“用大规模并行隐藏延迟”。因此,单个线程的执行速度(延迟)往往不是首要优化目标,而整体吞吐量(每秒完成的指令或事务数)才是关键。这四个指标分别从不同层面刻画系统状态:

  • 吞吐量:单位时间内完成的操作数,常用单位有 GFLOPs(浮点运算)、GT/s(纹理采样)、GIPS(指令)。它反映计算资源或特定硬件单元(如 FP32 核心、Tensor Core)的利用程度。
  • 延迟:单个操作从发出到完成的时间间隔,通常以时钟周期或微秒计。在 GPU 中,内存访问延迟(数百周期)远高于算术指令延迟(几个周期)。延迟本身难以直接优化,但可以通过提高并行度来“隐藏”。
  • 占用率:每个流多处理器(SM)上活跃线程束(Warp)数量与最大可容纳线程束数量的比值。高占用率是隐藏延迟的前提,但并非越高越好——过高的占用率可能导致寄存器溢出或缓存抖动。
  • 内存带宽利用率:实际访问的字节数与理论峰值带宽的比值。它不等于“带宽占用率”,而是衡量访存请求是否高效利用了 DRAM 的突发传输特性。

下图展示了四个指标在典型 GPU 执行流水线中的位置与反馈关系:

指令流 Warp 调度器 流多处理器 (SM) 占用率 活跃 Warp / 最大 Warp 计算吞吐量 FP32 / INT / Tensor 延迟隐藏机制 高占用率 → 更多 Warp 可切换 → 隐藏访存延迟 内存层次与带宽 L1 / L2 缓存 命中率影响延迟 DRAM 带宽 利用率 = 实际/峰值 带宽利用率 % of peak 吞吐量 = 计算单元利用率 × 时钟频率;延迟 = 访存路径总耗时;两者通过占用率耦合

图中可见,占用率是连接指令流与计算/访存单元的枢纽。当占用率不足时,SM 无法用其他 Warp 掩盖当前 Warp 的访存等待,延迟就会直接暴露为流水线空泡,拉低整体吞吐量。而内存带宽利用率则取决于访存模式是否连续、是否充分利用缓存行。

3.2 方法流程:Nsight 中如何采集与解读指标

使用 Nsight Compute(ncu)进行指标采集时,建议遵循以下流程:

  1. 粗筛:先运行 ncu --set full 获取全量指标,观察是否存在明显的“红绿灯”警告(如低占用率、高延迟)。
  2. 聚焦:根据粗筛结果,使用 --metrics 参数定向采集。例如:
    ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
    gpu__time_duration.sum,\
    sm__warps_active.avg.pct_of_peak_sustained_active \
    ./my_kernel
  3. 关联分析:将吞吐量与占用率数据绘制成时间线,观察是否存在“占用率低导致吞吐量峰值无法达到”的时段。
  4. 访存剖析:使用 l1tex__throughput.avg.pct_of_peak_sustained_elapsed 和 dram__throughput.avg.pct_of_peak_sustained_elapsed 对比 L1/L2 与 DRAM 的瓶颈。

下表总结了四个指标在 Nsight Compute 中的典型对应计数器及优化方向:

指标Nsight Compute 关键计数器理想范围优化方向
吞吐量sm__throughput.avg.pct_of_peak_sustained_elapsed> 80%减少指令依赖,使用向量化指令
延迟gpu__time_duration.sum + 内存延迟采样越低越好提高缓存命中,预取数据
占用率sm__warps_active.avg.pct_of_peak_sustained_active50%~75%调整 block 大小,减少寄存器使用
带宽利用率dram__throughput.avg.pct_of_peak_sustained_elapsed> 60%合并访问,使用只读缓存

3.3 案例:矩阵乘法中的指标权衡

以 C = A × B(尺寸 4096×4096)为例,使用朴素三重循环实现时,Nsight 报告显示:DRAM 带宽利用率仅 35%,而计算吞吐量达到 92%。这说明内核并非带宽瓶颈,而是计算瓶颈。但进一步观察占用率发现,由于每个线程需要 64 个寄存器,占用率只有 37%,导致延迟隐藏不足。

改用共享内存分块(tile size = 32×32)后,寄存器压力降至 40 个,占用率提升至 62%。此时 DRAM 利用率上升至 58%,计算吞吐量略降至 85%,但总执行时间缩短了 40%。这个案例揭示了一个关键原则:不要单独追求某个指标最大化,而要寻找平衡点。高占用率不一定带来高性能——如果每个线程的局部性变差,缓存命中率下降,反而增加 DRAM 压力。

3.4 误区与复盘:常见指标解读陷阱

误区一:将“低占用率”直接等同于“优化失败”。某些内核(如归约操作)天然需要低占用率来避免原子操作冲突。此时应关注是否有其他 Warp 可切换,而非盲目增加 block 数量。

误区二:忽略“延迟”与“吞吐量”的因果关系。当 DRAM 利用率高但计算吞吐量低时,问题往往不是带宽不足,而是延迟过长导致 SM 等待数据。此时应检查 L2 命中率,而非直接增加并行度。

误区三:只看平均带宽利用率,忽略访问模式。两个内核可能都报告 70% 的 DRAM 利用率,但一个为连续访问,另一个为随机访问。随机访问会导致 DRAM 行切换频繁,实际有效带宽远低于理论值。Nsight 的 lts__t_sector_hit_rate 可辅助判断。

复盘建议:每次优化后,同时记录四个指标的变化,并绘制“指标雷达图”。若某项指标提升但总时间未缩短,说明瓶颈已转移至其他资源。例如,提高占用率后计算吞吐量下降,可能因为 L1 缓存抖动加剧。此时应回退优化,或改用 __launch_bounds__ 限制最大线程数,以保留更多 L1 缓存。

3.5 小结与实操提示

本章的四个指标构成了 GPU 性能分析的“四维坐标系”。在后续章节使用 Nsight Systems 查看 CPU-GPU 整体时间线时,请始终带着这些问题:当前内核的占用率是否足以隐藏访存延迟?DRAM 利用率是否接近峰值?计算单元是否空闲等待数据?只有将指标与内核代码结构对应起来,才能从“看到数字”进阶到“读懂行为”。

实操建议:在 Nsight Compute 的“Details”视图中,点击“Memory Workload Analysis”标签页,查看 Mem Busy% 与 Max Bandwidth% 的差值。若两者差距大于 20%,说明存在严重的访存延迟,而非带宽饱和。此时优先优化数据局部性,而非增加并发。