内核分析基础：如何启动内核分析、收集性能指标、理解分析报告

核心原理：从“跑一次”到“看透一次”

在CUDA程序开发中，我们通常先关注功能正确性，再关注性能。但“程序能跑”和“程序跑得快”之间隔着巨大的鸿沟。Nsight Compute（以下简称Ncu）的核心价值在于：它不只是告诉你内核耗时多少，而是深入GPU内部，揭示SM（流多处理器）级的执行细节——哪些流水线在等待、哪些资源成为瓶颈、哪些指令产生了不必要的开销。

理解Ncu的分析逻辑，需要先建立三个基本认知：

  • 基于硬件计数器的采样：Ncu通过GPU硬件性能计数器（Performance Counter）采集数据，而非简单的时间戳。这意味着它能精确到指令周期、缓存命中率、访存事务数等底层指标。
  • 瓶颈树（Speed of Light）模型：Ncu将内核性能归约为“计算吞吐”“访存带宽”“延迟隐藏”三大维度，并以百分比形式展示理论峰值利用率。这能帮你快速判断内核属于计算密集型还是访存密集型。
  • 多级分析视图：从“概要（Summary）”到“详细页（Detail Pages）”，再到“源代码关联（Source View）”，Ncu逐层下钻，让性能问题从“现象”定位到“指令”再到“代码行”。
关键认知：Ncu的分析结果不是“评分”，而是“证据链”。每个性能问题（如访存延迟过高）都会关联到具体的指令、内存访问模式或资源占用情况。学会阅读这条证据链，比记住单个指标数值更重要。

方法流程：启动分析、收集指标、导出报告

启动Ncu分析通常有两种方式：命令行（CLI）和图形界面（GUI）。在服务器或自动化脚本场景中，CLI是首选；在交互式调优场景中，GUI的“向导式”分析更直观。下面以CLI为例，展示标准流程。

第一步：基础分析命令

# 对可执行文件中的某个内核进行默认分析
ncu --kernel-name "myKernel" --launch-count 1 ./myApp

# 如果程序有多个内核，且不确定名称，可先用 --list 查看
ncu --list ./myApp

默认分析会收集一组“基础指标集”，包括执行时间、SM利用率、访存吞吐等。但若要深入定位问题，需要指定更详细的指标集。

第二步：收集关键性能指标

Ncu将指标分为多个集合，常用组合如下：

指标集名称覆盖内容适用场景
SpeedOfLight计算/访存吞吐利用率快速判断瓶颈类型
MemoryWorkloadAnalysis缓存命中率、全局/共享访存事务访存密集型内核
ComputeWorkloadAnalysis指令混合、流水线利用率计算密集型内核
LaunchStats线程块/线程束配置、占用率资源占用与延迟隐藏分析
WarpStateStats线程束停滞原因分布定位延迟来源（如访存等待）
# 收集访存与线程束状态指标
ncu --set detailed --kernel-name "myKernel" ./myApp

# 自定义指标：只收集特定计数器
ncu --metrics gpu__time_duration.sum,sm__throughput.avg.pct_of_peak_sustained_elapsed,smsp__warp_issue_stalled_long_scoreboard_per_warp_active.pct ./myApp
注意：不要一次性收集所有指标。Ncu的计数器复用机制会导致同一计数器被多次重放（Replay），显著增加分析时间。建议先跑一次--set basic，再根据初步结果选择--set detailed或自定义指标集。

第三步：生成并导出分析报告

# 导出为文本格式（适合脚本解析）
ncu --export report.ncu-rep --force ./myApp

# 导出为CSV（适合导入Excel或pandas）
ncu --export report.csv --csv --kernel-name "myKernel" ./myApp

GUI模式下，点击“Export”按钮同样可以生成.ncu-rep文件，该文件可在Ncu GUI中重新打开，便于后续复盘或团队共享。

案例场景：一个访存瓶颈内核的完整分析

假设我们有一个简单的向量求和内核，但性能远低于预期。代码如下：

__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];  // 每个线程访问3个全局数组
    }
}

启动Ncu分析并查看SpeedOfLight结果：

ncu --set speedoflight --kernel-name "vectorAdd" ./vecAdd

输出关键片段如下：

Memory Throughput: 78.4%
Compute Throughput: 12.6%
L1/TEX Cache Hit Rate: 62.3%
L2 Cache Hit Rate: 95.1%

从数据可以判断：访存吞吐接近饱和（78.4%），而计算吞吐很低（12.6%），这是一个典型的访存密集型内核。进一步查看MemoryWorkloadAnalysis，发现全局加载请求的扇出（Sector）效率只有约70%，意味着部分缓存行没有被完全利用。

结合WarpStateStats，发现线程束停滞原因中Long Scoreboard占比高达65%，说明线程在等待全局内存数据返回。此时可以得出结论：瓶颈在于全局访存延迟未能被充分隐藏，且存在一定的访存事务浪费。

优化方向自然浮现：使用__ldg()或只读缓存、调整数据布局以提高空间局部性、增加每个线程处理的数据量（即网格跨步循环）以减少访存指令总数。

复盘要点：该案例中，Ncu的价值在于将“程序慢”这一模糊现象，拆解为“访存吞吐高、扇出效率低、长延迟等待占比高”三个可验证的事实。每个事实都指向具体的代码修改策略。

误区与复盘：新手常犯的五个分析错误

  • 只看执行时间，不看利用率：时间短不代表效率高，可能只是问题规模小。应结合SpeedOfLight百分比判断资源是否被充分利用。
  • 忽略“重放”开销：收集过多指标导致Ncu多次重放内核，分析时间成倍增加。应分层收集，先粗后细。
  • 混淆“瓶颈”与“现象”：例如看到Long Scoreboard高，就认为是“访存慢”，但真正原因可能是线程束数量不足，无法用其他线程束掩盖延迟。需要结合LaunchStats中的占用率一起看。
  • 不检查指标的有效性：某些计数器在特定GPU架构上不支持或返回0。分析前应确认目标架构（如Hopper、Ada Lovelace）与Ncu版本兼容。
  • 只分析一次运行：GPU频率波动、后台任务干扰可能导致单次结果偏差。建议至少分析3次，取中位数或平均值。

下面这张图总结了从启动分析到定位瓶颈的完整决策路径，帮助你建立系统性的分析思维：

Ncu内核分析决策路径 ① 启动分析 ncu --set basic ② 初步判断瓶颈 SpeedOfLight 计算 vs 访存 访存密集? ③a 访存深度分析 MemoryWorkloadAnalysis L1/L2命中率、扇出效率 ③b 计算深度分析 ComputeWorkloadAnalysis 指令混合、流水线利用率 ④ 延迟与占用率分析 WarpStateStats + LaunchStats 停滞原因 vs 线程束数量 定位根因 → 制定优化策略 例如：提高扇出效率 / 增加线程束 / 使用只读缓存 / 调整数据布局 常见误区 • 只看时间不看利用率 • 忽略计数器重放开销 • 混淆瓶颈与现象 • 单次运行即下结论 最佳实践 • 分层收集指标 • 结合多个视图交叉验证 • 至少3次取中位数

上图展示了从启动分析到定位根因的完整闭环。值得注意的是，访存密集与计算密集的判定并非绝对，许多内核处于两者之间。此时应优先分析占比更高的一方，再考虑另一方的优化空间。

最后，复盘时建议将Ncu报告与代码版本关联保存。例如，在Git提交信息中附带ncu --export生成的报告文件，或使用ncu --csv输出关键指标并纳入CI/CD流水线。这样，每次代码变更都能自动追踪性能变化，避免“优化后又回退”的尴尬情况。

← 上一章 📖 返回目录 下一章 →