3、NVIDIA Nsight Systems入门：Nsight Systems安装、时间线分析、CPU/GPU活动追踪
好，咱们今天聊聊Nsight Systems。说实话，这工具是我日常调试GPU性能问题时，打开频率最高的一个。没有之一。
很多刚入行的同学喜欢一上来就怼Nsight Compute，盯着kernel的occupancy和带宽看半天。我个人的习惯是——先别急。先拿Nsight Systems拉一条时间线，看看你的CPU和GPU到底在干嘛。很多时候，性能瓶颈根本不在kernel内部，而是CPU在等GPU，或者GPU在等数据。
3.1 安装与配置
安装其实没什么好说的，去NVIDIA官网下载就行。但我提醒一句：版本一定要匹配。
我曾经遇到过一回，项目组新来的同事装了个最新版的Nsight Systems，结果连不上老驱动。折腾了两天，最后发现是驱动版本太旧，Nsight Systems的新特性不兼容。嗯，这里要注意：
- Windows：直接下载.exe安装包，一路Next就行
- Linux：推荐用.run文件，或者通过apt/yum添加NVIDIA的仓库
- 容器环境：如果你在Docker里跑，记得挂载
/tmp/.X11-unix和/dev/shm
nsys-ui 打开图形界面，nsys profile 做命令行采集。这样方便集成到自动化脚本里。
3.2 第一次采集：从命令行开始
图形界面虽然直观，但我建议你先学会命令行。为什么？因为生产环境往往没有显示器。
最简单的采集命令长这样：
nsys profile -o my_profile ./my_app这行命令会运行你的程序，并在结束后生成一个my_profile.qdrep文件。双击就能在Nsight Systems里打开。
如果你想看得更细一点，可以加几个参数：
nsys profile --trace=cuda,nvtx,osrt -o my_profile ./my_app--trace=cuda：追踪CUDA API调用和kernel启动--trace=nvtx：追踪你自己插入的NVTX标记（后面会讲）--trace=osrt：追踪操作系统运行时，比如线程调度、内存分配
你想想看，这三个选项一开，基本上CPU和GPU之间的交互就一览无余了。
3.3 时间线分析：看懂这张图
打开生成的.qdrep文件，你会看到一张时间线图。别慌，我来带你拆解。
时间线从上到下分几个区域：
| 区域 | 显示内容 | 我关注什么 |
|---|---|---|
| CPU Threads | 每个CPU线程的活动 | 有没有线程在空转？有没有锁竞争？ |
| CUDA HW | GPU硬件引擎的活动 | 计算和拷贝是否重叠？ |
| CUDA API | CPU调用的CUDA API | cudaMemcpy是不是同步的？ |
| NVTX | 用户自定义标记 | 代码逻辑是否按预期执行？ |
我一般先看CUDA HW这一行。如果看到一大片空白，说明GPU在等CPU喂数据。如果看到计算和拷贝串行执行，说明没有用stream做流水线。
3.4 CPU/GPU活动追踪：找到真正的瓶颈
光看时间线还不够，你得学会定位问题。我总结了三类最常见的性能杀手：
3.4.1 CPU端瓶颈
如果你看到CPU线程长时间处于绿色（运行态），但GPU那边是空的——说明CPU在忙别的事，没来得及给GPU发任务。
我曾经调试过一个图像处理管线，发现CPU在预处理阶段花了大量时间做内存拷贝。后来改成用cudaMallocHost分配锁页内存，CPU端的拷贝时间直接砍半。
3.4.2 GPU端瓶颈
如果GPU的计算引擎和拷贝引擎没有重叠，说明你的stream用错了。或者，你根本就没用stream。
举个例子：
// 错误做法：串行执行
cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(dst);
cudaMemcpy(result, dst, size, cudaMemcpyDeviceToHost);
// 正确做法：用stream重叠
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream1);
kernel<<<grid, block, 0, stream2>>>(dst);
cudaMemcpyAsync(result, dst, size, cudaMemcpyDeviceToHost, stream3);在Nsight Systems的时间线上，你能清楚地看到三个stream的活动是否重叠。如果不重叠，那就是你的stream同步点没设对。
3.4.3 同步开销
这个最隐蔽。你可能会看到CPU和GPU都在忙，但吞吐量就是上不去。这时候要检查cudaDeviceSynchronize的调用频率。
我见过最夸张的一个案例，有人在每个kernel后面都加了一个cudaDeviceSynchronize。结果GPU利用率只有30%，因为CPU一直在等同步完成。去掉多余的同步点之后，吞吐量直接翻了2倍。
cudaDeviceSynchronize是性能杀手。除非你确实需要等待所有GPU操作完成（比如在计时、或者准备读取结果时），否则不要轻易调用。
3.5 用NVTX给时间线加注释
Nsight Systems默认只能看到CUDA API和kernel的名字。但你想啊，如果你的代码有几十个kernel，光看名字根本不知道哪个对应哪个逻辑。
这时候NVTX就派上用场了。你可以在代码里插入标记：
#include <nvtx3/nvToolsExt.h>
// 标记一个范围的开始
nvtxRangePushA("Preprocessing");
// ... 你的预处理代码 ...
// 标记结束
nvtxRangePopA();在Nsight Systems的时间线上，这些标记会以彩色条带的形式显示出来。你可以一眼看出哪个阶段耗时最长。
我个人习惯在每个关键函数入口和出口都加上NVTX标记。调试的时候，配合时间线一看，哪个函数拖了后腿，一目了然。
3.6 实战小技巧
- 别一次采集太久：Nsight Systems默认会记录所有事件。如果你的程序跑几个小时，生成的
.qdrep文件会大到离谱。我一般只采集关键几秒的数据。 - 用
--duration参数控制采集时长：nsys profile --duration=10 -o my_profile ./my_app只采集前10秒。 - 对比两次运行：优化前跑一次，优化后跑一次，把两个
.qdrep文件拖到一起对比。这是我最常用的方法。
好了，Nsight Systems的入门就讲到这里。说白了，这工具就是给你一双透视眼，让你看清CPU和GPU之间到底在干什么。下一章我们会深入Nsight Compute，专门分析kernel内部的性能问题。