汤不热吧NVIDIA Nsight Systems (NSS) 是一个强大的系统级性能分析工具，它能够帮助开发者深入了解 GPU 和 CPU 之间的交互，以及 CUDA 算子（Kernel）的执行效率。对于深度学习应用而言，理解哪个算子耗时最长是性能优化的第一步。
本文将手把手教你如何收集数据并解读 NSS 报告，专注于定位算子瓶颈。
步骤一：准备与数据采集
要使用 Nsight Systems，你需要安装 NVIDIA 驱动和 Nsight Systems 工具包。数据采集通常通过命令行接口（CLI）进行，因为它对应用性能的影响最小。
假设我们有一个名为 my_dl_model.py 的深度学习脚本。
采集命令示例
我们指定采集 CUDA API 调用 (cuda) 和 NVTX（用于用户自定义标记，如果有的话），并要求输出统计数据 (–stats true)：
# 确保你的环境路径包含 nsys
# -t cuda: 采集 CUDA 活动
# --stats true: 在控制台打印出汇总统计信息
# -o report_name: 指定输出文件名称
nsys profile -t cuda --stats true -o my_gpu_profile python my_dl_model.py
执行完毕后，你会得到一个 .qdrep 文件（如 my_gpu_profile.qdrep），这是我们分析的基础数据。
步骤二：加载与初始视图解析
打开 Nsight Systems UI (Windows/Linux/MacOS 均可)，然后加载 .qdrep 文件。
1. 概览视图（Timeline View）
这是 NSS 最重要的视图。它将时间轴分为 CPU/Host 线程、CUDA API 调用、以及 GPU 上的活动（Kernel）。
关键观察点：
- Host Threads: 查看哪个 CPU 线程负责启动 CUDA 任务（通常是主线程或特定的 Worker 线程）。
- CUDA API: 查看 cudaLaunchKernel 或 cudaMemcpy 等 API 调用，这是 Host 向 Device 发送指令的记录。
- CUDA Streams: 这是 GPU 上的执行队列。每个 Kernel 都属于一个 Stream。我们的重点将集中在这些 Stream 上。
2. 定位瓶颈：聚焦 CUDA Kernels
将 Timeline 视图放大（Zoom In）到你感兴趣的执行阶段（例如，一个训练迭代或一次推理）。你会看到一系列颜色块，这些就是 GPU 上运行的 CUDA Kernel。
如何识别瓶颈 Kernel：
耗时的 Kernel 表现为 Stream 中占据最长时间的矩形块。它们通常是矩阵乘法（如 volta_sgemm 或 turing_hgemm）、卷积或特定的元素级操作（如 elementwise_add）。
点击任何一个 Kernel 块，底部的 Details 窗格会显示详细信息，包括：
- Duration: 实际执行时间。
- Grid/Block Dimensions: 启动配置。
- Source: 如果可能，显示其所属的 CUDA API 调用。
步骤三：利用统计视图（Statistics）快速定位
如果你的应用程序运行时间很长，通过 Timeline 一条条查找效率太低。这时，Statistics 视图是你的救星。
在左侧导航栏中选择 Statistics，然后导航到 GPU Activities 或 CUDA Kernels 标签页。
查找“Top Kernels by Duration”
NSS 会自动聚合所有相同名称的 Kernel 的执行时间。通常，表格会按总耗时（Total Duration）降序排列。
| Kernel Name | Total Duration (ms) | % of Total Time | Average Duration (ms) |
|---|---|---|---|
| volta_sgemm_128x128_nn | 550.8 | 45.1% | 2.5 |
| cudnn::convolution::compute_hgrad | 210.1 | 17.2% | 1.1 |
| cudaMemcpyAsync | 90.5 | 7.4% | 0.05 |
解读：
上例清晰地表明，名为 volta_sgemm_128x128_nn 的 GEMM（通用矩阵乘法）Kernel 占据了总 GPU 运行时间的近一半。这意味着这是最主要的算子瓶颈，优化工作应优先聚焦于它（例如，通过模型量化、使用更高效的 Tensor 布局或调整 Batch Size）。
步骤四：识别 Host/Device 间隙（Serialization Gaps）
定位算子瓶颈的“艺术”在于区分“算子本身运行慢”和“算子之间有不必要的等待”。
回到 Timeline 视图，观察同一 CUDA Stream 中相邻 Kernel 之间的空隙（Gaps）。
- 如果 Kernel 本身很长，且 Stream 之间没有空隙： 这是纯粹的计算瓶颈，需要优化该 Kernel 的效率。
- 如果 Kernel 很短，但它们之间有明显的空隙： 这通常是 Host（CPU）端启动下一个 Kernel 的延迟（Host Overhead），也称为序列化瓶颈。这可能由以下原因造成：
- 同步点过多（如不必要的 cudaDeviceSynchronize()）。
- CPU 端数据预处理耗时过长，阻塞了下一个 Kernel 的启动。
通过 NSS，你可以将 Timeline 中的空隙与上方 Host 线程中的 CUDA API 调用进行关联，从而确认是哪一个 API 调用导致了延迟，最终实现精准的性能优化。