4、NVIDIA Nsight Compute核心功能：内核分析器、性能指标解读、瓶颈定位
好，咱们今天来聊聊 Nsight Compute。说实话，这工具是我日常工作中用得最多的一个。你想想看，写一个 CUDA 内核，跑起来发现性能不对，你总得知道问题出在哪吧？Nsight Compute 就是干这个的。它不像 nvprof 那样只给你看个大概，它能深入到每个 warp、每条指令，告诉你到底哪里在“卡脖子”。
4.1 内核分析器：不只是“跑一下”那么简单
很多人用 Nsight Compute，就是点一下“Profile”，然后看结果。嗯，这当然可以，但我觉得有点浪费。我个人习惯是，先想清楚我要分析什么，再动手。
内核分析器其实提供了两种模式：Profile 和 Trace。Profile 模式会收集内核的详细性能计数器，比如指令吞吐、内存带宽、占用率等等。Trace 模式则更关注时间线，能看到每个 kernel launch 的耗时、API 调用顺序。
我建议你刚开始做性能分析时，先用 Profile 模式。它能给你一张“体检报告”，告诉你哪里可能有问题。比如，我曾经遇到一个内核，跑得很慢，直觉告诉我可能是访存问题。结果 Profile 一跑，发现 L1 缓存命中率只有 30%。嗯，问题找到了。
核心要点：Profile 模式看“是什么”，Trace 模式看“什么时候”。先 Profile，再 Trace，这是标准流程。
4.2 性能指标解读：别被数字骗了
Nsight Compute 会给你一大堆指标，从占用率到带宽利用率，从指令 mix 到 stall 原因。说实话，新手很容易看花眼。我刚开始接触时也这样，看到一堆数字，不知道哪个才是关键。
我个人习惯是，先看三个核心指标：
- 占用率（Occupancy）：每个 SM 上活跃的 warp 数量。占用率低，说明你的内核没有充分利用硬件资源。但注意，占用率不是越高越好。我记得有一次，我为了追求高占用率，把寄存器用得很少，结果导致频繁的本地内存溢出，性能反而下降了。
- 计算吞吐（Compute Throughput）：看你的内核是计算密集型还是访存密集型。如果计算吞吐接近理论峰值，那问题可能出在访存上。
- 访存吞吐（Memory Throughput）：看你的内核是否把显存带宽用满了。如果带宽利用率很低，说明你的访存模式有问题，比如非合并访问。
举个例子，你看下面这个表格，是我从一个实际项目中摘出来的：
| 指标 | 值 | 说明 |
|---|---|---|
| Occupancy | 75% | 还不错，但还有提升空间 |
| Compute Throughput | 40% | 计算单元没跑满，可能是访存瓶颈 |
| Memory Throughput | 85% | 带宽利用率很高，说明访存压力大 |
| L1 Hit Rate | 45% | 偏低，需要优化数据局部性 |
你看，这个内核的占用率 75%，看起来还行。但计算吞吐只有 40%，而访存吞吐高达 85%。这说明什么？说明内核被访存“拖死”了。你再看看 L1 命中率只有 45%，嗯，问题就在这。
小技巧：Nsight Compute 里有个“Speed of Light”图，能直观地看到你的内核离理论极限还有多远。我每次优化完，都会看一眼这个图，心里就有底了。
4.3 瓶颈定位：从指标到代码
指标看完了，知道问题在哪了，接下来就是定位到具体的代码行。Nsight Compute 的 Source View 功能就是干这个的。它能把你内核的每一行代码和对应的性能计数器关联起来。
比如，你发现某个循环的访存延迟很高，Source View 会高亮显示那几行代码，并告诉你具体的 stall 原因。是“等待数据从全局内存返回”？还是“等待计算单元空闲”？
我举个例子，假设你有这样一段代码：
__global__ void myKernel(float* A, float* B, float* C, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
// 这行访存可能有问题
float a = A[idx];
float b = B[idx];
// 这行计算
C[idx] = a + b;
}
}在 Nsight Compute 的 Source View 里，你会看到 float a = A[idx]; 这行旁边有个红色的标记，表示这里 stall 很高。点进去一看，stall 原因是 Long Scoreboard，说白了就是等待全局内存数据返回。
为什么会这样？因为 A[idx] 和 B[idx] 是两次独立的全局内存访问，而且没有合并。你想想看，每个线程都去读自己的 idx，如果 idx 是连续的，那没问题，硬件会自动合并成一次大的访存。但如果 idx 是随机的，那就惨了，每次访存都是小包，带宽利用率极低。
我曾经在一个项目中遇到过类似的问题。一个矩阵转置的内核，性能死活上不去。用 Nsight Compute 一查，发现访存模式是“非合并”的。后来我把共享内存用上，先让线程协作把数据搬到共享内存，再转置写回，性能直接翻了 3 倍。
避坑指南：我曾经犯过一个错误，就是只看全局指标，不看 Source View。结果优化了半天，发现改的地方根本不是瓶颈。记住，指标告诉你“是什么”，Source View 告诉你“在哪里”。两者缺一不可。
4.4 实战：一个完整的分析流程
好了，理论讲完了，咱们来个实战。假设你有一个内核，性能不达标。我建议你按这个流程走：
- 跑一次 Profile：看看整体指标，判断是计算瓶颈还是访存瓶颈。
- 看 Speed of Light 图：直观感受一下离理论极限有多远。
- 看 Stall 原因：在“Warp State”里，看看 warp 到底在等什么。是等数据？等计算？还是等同步？
- 打开 Source View：找到具体的代码行，确认问题。
- 修改代码：根据分析结果，优化访存模式、减少分支、提高占用率等。
- 重新 Profile：验证优化效果。如果没达到预期，回到第 1 步。
嗯，这个流程我用了好几年，基本没出过问题。你刚开始可能觉得麻烦，但用多了就会发现，这其实是最快的方式。
总结一下：Nsight Compute 的核心价值，就是帮你把“感觉”变成“数据”。别靠猜，让工具说话。记住，性能优化不是玄学，是科学。
好了，这一章就到这里。下一章咱们聊聊怎么用 Nsight Compute 分析更复杂的内核，比如那些用了动态并行和 Cooperative Groups 的。到时候再给你分享一些我踩过的坑。