4、NVIDIA Nsight Compute入门：Kernel Profiling、Metrics分析、Source Correlations

好，咱们今天聊聊 Nsight Compute。说实话，这工具刚出来那会儿，我还有点抵触——毕竟 Nsight Visual Studio Edition 用习惯了。但后来被一个性能问题逼得没办法，硬着头皮试了试，结果真香。现在它已经是我做 GPU 性能分析的标配了。

4.1 为什么是 Nsight Compute？

你可能会问：用 nvprof 或者 ncu 命令行不就行了？嗯，命令行确实能拿到数据，但 Nsight Compute 的 GUI 版本，说白了就是把那些枯燥的数字变成了可视化的图表和关联代码。我个人习惯是先用命令行跑一轮，拿到大概的瓶颈方向，然后打开 GUI 做精细分析。

Nsight Compute 的核心能力就三个：

• Kernel Profiling：对单个 kernel 做深度剖析，不是全局那种粗粒度统计
• Metrics 分析：几百个硬件计数器，帮你定位到底是计算慢还是访存慢
• Source Correlations：把性能数据直接映射到你的 CUDA 代码行上

我在项目中遇到过最典型的一个场景：一个 kernel 跑得很慢，但用 nvprof 看 occupancy 很高，SM 利用率也高，就是不知道问题在哪。后来用 Nsight Compute 打开 Source Correlations，发现有一行代码的 stall 原因全是 "long scoreboard"——说白了就是访存等待。再点进去看，原来是一个全局内存的随机访问模式没处理好。

4.2 快速上手：Profiling 一个 Kernel

咱们先走一遍流程。假设你有一个编译好的可执行文件，里面调用了你的 CUDA kernel。

第一步：启动 Nsight Compute GUI

在终端里敲 ncu-ui，或者从开始菜单启动。我个人建议用命令行启动，因为可以带一些参数，比如指定 GPU 序号。

第二步：创建一个 Profile 会话

点击 "Create a new session"，然后指定你的可执行文件路径和参数。这里有个小坑：如果你用相对路径，最好把工作目录也设对，不然可能找不到依赖的库文件。

第三步：运行并收集数据

点击 "Launch" 按钮。Nsight Compute 会启动你的程序，并在每个 CUDA kernel 执行时收集详细的性能计数器数据。注意，它默认只收集第一个 kernel 的完整数据，如果你想看所有 kernel，需要在设置里改一下。

4.3 Metrics 分析：看懂那些数字

数据收集完了，你会看到一堆指标。别慌，咱们挑几个最重要的说。

我记得有一次，一个 kernel 的 occupancy 高达 90%，但 Duration 还是很大。我盯着 Metrics 看了半天，发现 "Memory Throughput" 只有理论峰值的 15%。嗯，这就说明问题不在 occupancy，而在访存带宽没利用起来。后来发现是每个线程只读了一个 float，但 stride 很大，导致 cache line 浪费严重。

4.4 Source Correlations：把问题定位到代码行

这个功能是我最喜欢 Nsight Compute 的地方。它能把性能数据直接映射到你的 CUDA 源代码上。

操作很简单：在 "Source" 标签页里，你会看到你的 kernel 代码，每一行前面都有一个颜色条。颜色越红，说明这一行消耗的时间越多。点击某一行，右侧会显示详细的 stall 原因分析。

举个例子，假设你有这样一段代码：

__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        // 这一行可能很慢
        C[idx] = A[idx] + B[idx];
    }
}

如果 Source Correlations 显示 C[idx] = A[idx] + B[idx] 这一行是红色的，点进去看 stall 原因，发现 "long scoreboard" 占了 80%。这说明什么？说明访存延迟是瓶颈。为什么？因为每个线程访问的地址是连续的，但 GPU 的 memory coalescing 可能没做好——比如你的数组是 float 类型，但线程束的访问模式不对齐。

我曾经遇到过一个更隐蔽的问题：代码里有一行 if (data[idx] > threshold)，Source Correlations 显示这行很慢。我一开始以为是分支发散，但仔细看 stall 原因，发现是 "instruction fetch" 等待。后来才意识到，是因为这行代码所在的 basic block 太大，指令缓存不够用了。解决办法是把条件判断提前，减少指令数。

4.5 实战：一个完整的分析流程

好了，理论说完了，咱们走一个完整的例子。假设你有一个矩阵乘法的 kernel，性能不理想。

1. 第一步：用 ncu 命令行跑一次

   ncu --set full -o matmul_profile ./matmul

   这个命令会收集所有可用的 metrics，并保存到 matmul_profile.ncu-rep 文件里。

2. 第二步：用 ncu-ui 打开结果

   双击 .ncu-rep 文件，或者用 ncu-ui matmul_profile.ncu-rep 打开。

3. 第三步：看 "Summary" 页面

   先看 Duration 和 Occupancy。如果 occupancy 低，检查是不是寄存器用太多或者 shared memory 太大。

4. 第四步：看 "Memory Workload Analysis"

   如果 occupancy 没问题，但 Duration 大，看 Memory Throughput。如果吞吐量低，点进去看是 L1、L2 还是 HBM 的瓶颈。

5. 第五步：看 "Source" 页面

   找到最红的代码行，分析 stall 原因。如果是 "long scoreboard"，优化访存模式；如果是 "no instruction"，优化指令并行度。

我记得有一次，我用这个流程分析一个卷积 kernel。发现 occupancy 很高，但 Memory Throughput 只有 30%。Source Correlations 显示最慢的是 input[idx] * weight[j] 这一行，stall 原因是 "long scoreboard"。我仔细一看，原来 input 的访问模式是 input[threadIdx.x * stride + offset]，stride 不是 1，导致 coalescing 失效。改成 input[threadIdx.x + offset] 后，吞吐量直接翻倍。

4.6 避坑指南

最后，分享几个我踩过的坑：

• 不要一次看太多 metrics：Nsight Compute 能收集几百个指标，但你看不过来。我建议先看 Duration、Occupancy、Memory Throughput、Compute Utilization 这四个，有问题再深入。
• 注意 profiling 的开销：有些 metrics 需要重放 kernel 才能收集，这会增加 profiling 时间。如果 kernel 本身就跑得很慢，建议先用 --set basic 快速扫一遍。
• Source Correlations 的精度问题：对于很短的 kernel（比如 1-2 微秒），采样可能不够。这时候可以用 --replay-mode kernel 强制重放，但时间会变长。
• 我曾经犯过一个错误：在 profiling 时开了太多后台程序，导致 GPU 被其他进程占用，结果 metrics 显示 SM 利用率很低。后来才发现是 profiling 环境不干净。建议 profiling 前关掉其他 GPU 应用，或者用 CUDA_VISIBLE_DEVICES 指定 GPU。

嗯，Nsight Compute 的内容差不多就这些。说白了，它就是一个放大镜，帮你把 GPU 内部的细节看清楚。但工具再好，也得靠经验去解读。我建议你拿一个自己写的 kernel 跑一遍，看看 Source Correlations 里最红的那行代码是什么，然后想想为什么。多练几次，你就能一眼看出瓶颈在哪了。