GPU性能分析：CUDA Profiling基础，使用nvprof与Nsight Systems
说到GPU性能分析，我脑子里第一个蹦出来的画面，是几年前调一个深度学习训练框架的场景。那时候代码跑得慢，我盯着屏幕干瞪眼，完全不知道瓶颈在哪。后来硬着头皮上了nvprof，才第一次看清GPU到底在忙什么——嗯，那感觉就像给黑盒子开了个天窗。
今天咱们就聊聊CUDA Profiling的基础。说白了，就是搞清楚你的GPU代码到底慢在哪。我个人习惯把性能分析分成两个层次：先用nvprof快速定位热点，再用Nsight Systems深挖时间线。这两个工具，一个像听诊器，一个像CT机。
为什么需要GPU Profiling？
你想想看，写CUDA代码的时候，我们总觉得“并行嘛，肯定快”。但实际跑起来，可能有一半的线程在空转，或者数据在显存和内存之间来回倒腾。我遇到过最夸张的一次，一个kernel实际计算只占了10%的时间，剩下90%都在等数据搬运。
没有Profiling，你就是在盲人摸象。有了它，你才能回答三个核心问题：
- 计算密集度够不够？ 每个线程干活的比重有多大
- 访存模式好不好？ 有没有合并访问、有没有bank conflict
- 并行度足不足？ occupancy是不是被压得太低
核心观点： 性能优化的第一步，永远不是改代码，而是测量。没有数据，你连问题在哪都不知道。
nvprof：老牌工具，快速上手
nvprof是CUDA Toolkit自带的命令行工具。虽然NVIDIA现在主推Nsight系列，但nvprof在快速排查时依然好用。我至今还留着几个常用的命令模板，直接复制粘贴就能用。
基本用法
# 最简单的用法，统计所有kernel的耗时
nvprof ./my_cuda_app
# 只看某个特定的kernel
nvprof --kernels myKernel ./my_cuda_app
# 导出时间线到文件，方便后续分析
nvprof -o timeline.nvprof ./my_cuda_app跑完之后，你会看到类似这样的输出：
==20972== Profiling result:
Time(%) Time Calls Avg Min Max Name
60.23% 1.2345s 10 123.45ms 100.12ms 150.34ms vectorAdd_kernel
25.12% 0.5123s 2 256.15ms 240.10ms 270.20ms matrixMul_kernel
14.65% 0.2987s 5 59.74ms 50.23ms 70.12ms reduce_kernel看到这个表格，你第一反应应该是：哪个kernel最耗时？ 上面这个例子，vectorAdd_kernel占了60%的时间，那优化重点就是它。
我的小技巧： 用 nvprof --print-gpu-trace 可以看每个kernel的详细时间线，包括启动延迟、执行时间、内存传输时间。我曾经靠这个发现一个kernel每次启动都有5ms的延迟，后来发现是没预热CUDA context。
常用指标解读
| 指标 | 含义 | 理想值 |
|---|---|---|
| Occupancy | 每个SM上活跃的warp比例 | 50%以上 |
| Global Load Efficiency | 全局内存加载效率 | 接近100% |
| Branch Divergence | 分支分歧率 | 越低越好 |
| Memcpy Bandwidth | 内存拷贝带宽 | 接近理论峰值 |
举个例子，如果你看到Occupancy只有20%，那说明你的block配置有问题。我遇到过新手把block size设成32，结果每个SM只能跑一个block，浪费了大量计算资源。
Nsight Systems：时间线的艺术
nvprof能告诉你“哪个kernel慢”，但Nsight Systems能告诉你“为什么慢”。它把整个程序的执行过程画成时间线，CPU活动、GPU kernel、内存传输、API调用，全都一目了然。
安装与启动
Nsight Systems是独立工具，可以从NVIDIA官网下载。启动后，选择你的可执行文件，点“Run”就行。我个人习惯先跑一次全量分析，再根据结果缩小范围。
# 命令行方式也可以
nsys profile -o my_report ./my_cuda_app时间线怎么看？
打开生成的报告，你会看到类似这样的布局：
- 最上面： CPU活动，包括API调用（cudaMemcpy、cudaLaunchKernel等）
- 中间： GPU kernel执行，不同颜色代表不同kernel
- 最下面： 内存传输，包括Host到Device、Device到Host
我一般先看CPU和GPU之间有没有明显的空闲期。比如CPU在等GPU结果，或者GPU在等CPU喂数据。这种“串行等待”往往是性能杀手。
避坑指南： 我曾经遇到一个情况，时间线上GPU kernel之间有大段空白。查了半天，发现是CPU端在kernel启动后做了大量同步操作。解决方案是用CUDA Stream把多个kernel流水线化，让GPU一直有活干。
实战：用Profiling定位一个典型瓶颈
咱们来看一个真实场景。假设你写了一个向量加法，但性能很差。用nvprof跑一下：
nvprof --metrics gld_efficiency ./vector_add输出显示Global Load Efficiency只有30%。这说明什么？说明全局内存访问没有合并。你想想看，每个线程读自己的数据，但如果地址不连续，GPU就得多次访问内存。
解决方案很简单：确保相邻线程访问相邻地址。比如用 threadIdx.x 作为索引，而不是 blockIdx.x * blockDim.x + threadIdx.x 的某种乱序组合。
改完之后再跑一次，效率提升到95%。嗯，这就是Profiling的价值——让你用数据说话，而不是靠猜。
SVG：CUDA Profiling知识体系图
总结一下
GPU性能分析不是玄学，是科学。nvprof给你数据，Nsight Systems给你画面。两者结合，你就能像X光一样看透你的CUDA代码。我个人建议：先跑nvprof看全局，再用Nsight Systems深挖细节。别一上来就开Nsight，信息太多反而容易迷失。
最后提醒一句：Profiling本身也有开销。别在生产环境一直开着，那是给自己挖坑。嗯，今天就聊到这，下次咱们聊聊怎么用这些数据指导优化。