3、Nsight Systems入门：安装与配置、Timeline视图初探、捕获第一个Profile

好，咱们正式开始动手了。

前面两章聊了不少概念，什么GPU硬件架构、任务调度流水线，听着挺热闹。但说句实在话，没有工具辅助，你根本看不到这些调度到底是怎么跑的。就像修车，光看说明书不行，你得把引擎盖打开，接上诊断仪，看实时数据流。

Nsight Systems，就是CUDA程序的那台「诊断仪」。这一章，我带你把它装好，跑通，抓到第一个Profile数据。

3.1 安装与配置：别在这步卡住

安装本身不复杂，但有几个坑，我当年都踩过。

3.1.1 下载与版本选择

去NVIDIA官网下载Nsight Systems。我个人习惯直接装独立版，不依赖Visual Studio或Eclipse。这样最干净，也最灵活。

  • Windows：下载 .exe 安装包，一路Next就行。注意安装路径不要有中文。
  • Linux：下载 .run 文件。我建议用 sudo sh -c "echo 0 > /proc/sys/kernel/perf_event_paranoid" 临时放开性能计数器权限，否则有些硬件事件抓不到。
  • 版本匹配：Nsight Systems版本最好和你的CUDA Toolkit版本对应。比如CUDA 12.x，就用Nsight Systems 2024.x。混搭也能用，但偶尔会出些奇怪的小毛病。
注意： 我曾经在Ubuntu 20.04上装过最新版Nsight，结果发现内核版本太低，perf事件支持不全。后来老老实实降了一个小版本才搞定。所以，如果装完发现某些功能点不了，先检查一下系统内核版本。

3.1.2 环境变量配置

装完之后，最好把Nsight的bin目录加到PATH里。这样你可以在终端里直接敲 nsys 命令。

# Linux 示例，加到 ~/.bashrc 里
export PATH=/opt/nvidia/nsight-systems/2024.3/bin:$PATH

Windows用户更简单，安装程序会自动处理环境变量。如果没生效，手动加一下系统PATH就行。

3.1.3 验证安装

打开终端，敲下面这行命令：

nsys --version

如果看到类似 NVIDIA Nsight Systems version 2024.3.1 的输出，恭喜你，装好了。

小技巧： 我第一次装完后，直接跑了一个空程序测试，结果报错说找不到CUDA驱动。其实是因为我忘了设置 LD_LIBRARY_PATH。如果你也遇到类似问题，检查一下CUDA的lib路径有没有被正确加载。

3.2 Timeline视图初探：你的第一张「调度地图」

装好之后，咱们先别急着抓数据。先打开Nsight Systems，看看它的主界面长什么样。

说白了，整个工具的核心就一个视图——Timeline视图。它把CPU和GPU上发生的所有事情，按时间轴画出来。

3.2.1 界面布局

打开Nsight Systems，你会看到这样的布局：

  • 左侧：进程/线程列表。显示所有被追踪的进程和线程。
  • 中间：Timeline画布。这是主战场，所有事件都在这里展示。
  • 右侧：详情面板。点击某个事件，这里会显示它的详细信息。
  • 顶部：工具栏。控制抓取、回放、缩放等操作。

嗯，这里要注意：Timeline视图默认是按时间顺序从左到右排列的。CPU活动在上方，GPU活动在下方。你一眼就能看出CPU在等GPU，还是GPU在等CPU。

3.2.2 核心概念：Row与Event

Timeline里每一行（Row）代表一个执行单元。比如：

  • CPU Thread Row：显示某个CPU线程上的函数调用、API调用。
  • GPU Stream Row：显示某个CUDA Stream上的Kernel执行、Memcpy操作。
  • GPU HW Queue Row：显示硬件队列里的实际调度情况。

每个彩色的小方块就是一个Event。比如一个Kernel启动、一次cudaMemcpy、一次同步操作。颜色越深，通常表示耗时越长。

核心洞察： 我刚开始用Nsight时，总盯着Kernel的执行时间看。后来才发现，真正拖慢性能的，往往是那些不起眼的「小方块」——比如CPU和GPU之间的同步点、内存拷贝的间隙。Timeline视图的价值，就是让你看到这些「隐藏的等待」。

3.3 捕获第一个Profile：动手吧

理论说再多，不如跑一次。咱们写一个最简单的CUDA程序，然后用Nsight抓它的Profile数据。

3.3.1 准备一个测试程序

随便写个向量加法的程序就行。比如：

// vector_add.cu
#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1 << 20;
    size_t bytes = n * sizeof(float);

    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;

    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);

    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    vecAdd<<<256, 256>>>(d_a, d_b, d_c, n);

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    printf("Result: %f\n", h_c[0]);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

编译它：

nvcc -o vector_add vector_add.cu

3.3.2 用nsys命令行抓取

我个人习惯先用命令行抓，因为可以精确控制抓取范围。打开终端，执行：

nsys profile -o my_first_profile ./vector_add

这条命令的意思是：

  • nsys profile：启动性能分析器。
  • -o my_first_profile：输出文件名为 my_first_profile，默认生成 .nsys-rep 文件。
  • ./vector_add：要分析的目标程序。

运行结束后，你会看到类似这样的输出：

Processing events...
[1] Vector Add: 100% (2/2)
Generated: /home/user/my_first_profile.nsys-rep
提示： 如果你只想抓GPU活动，可以加 --gpu-metrics-device=0 参数。这样能减少CPU侧的开销，让Profile数据更干净。我调试GPU瓶颈时经常这么干。

3.3.3 在GUI中打开并查看

双击生成的 .nsys-rep 文件，或者直接在Nsight Systems里打开它。你会看到类似下面的Timeline：

Nsight Systems Timeline 视图示例 0ms 5ms 10ms CPU Thread (main) cudaMalloc cudaMemcpy H2D Kernel cudaMemcpy D2H GPU Stream (default) Memcpy H2D Kernel Memcpy D2H GPU HW Queue (SM) Copy Engine SM Active Copy Engine CPU等待GPU Kernel完成 图例 CPU API调用 内存拷贝 Kernel执行 GPU Memcpy GPU Kernel 注意：CPU上的cudaMemcpy是阻塞调用，会等待GPU完成。而Kernel启动是异步的，CPU会立即返回。

看到这张图了吗？这就是你抓到的第一个Profile。你能清楚地看到：

  • CPU先调了 cudaMalloc 和 cudaMemcpy（蓝色和红色方块）。
  • 然后GPU开始执行Kernel（橙色方块）。
  • 最后CPU等待GPU把结果拷回来（红色方块）。

你想想看，如果没有这个Timeline，你只能靠猜。现在一目了然：CPU在等GPU，还是GPU在等数据？

第一个Profile的收获： 我当年第一次看到自己的程序在Timeline上的样子，说实话挺震撼的。原来我以为Kernel是瞬间完成的，结果发现它占了大部分时间。更让我意外的是，cudaMemcpy的耗时比Kernel还长。从那以后，我养成了一个习惯：任何优化之前，先抓一个Profile看看。

3.4 本章小结

这一章，咱们把Nsight Systems装好了，跑通了，也看到了第一个Timeline视图。说白了，就是让你知道：

  • 怎么装、怎么配、怎么验证。
  • Timeline视图长什么样，Row和Event是什么意思。
  • 怎么用命令行抓一个Profile，然后在GUI里看。

嗯，到这里，你已经有了最基本的「诊断仪」。下一章，咱们会深入Timeline视图的细节，看看那些花花绿绿的方块到底代表什么，以及怎么从里面挖出性能瓶颈。

记住：工具只是手段，洞察才是目的。别急着优化，先看清楚你的程序到底在干什么。

课后练习： 试着用不同的block size（比如128、512）重新编译运行上面的向量加法程序，然后用nsys抓Profile。看看Kernel的执行时间有什么变化？Timeline上的形状有什么不同？这能帮你直观理解「线程块大小对性能的影响」。

想要更多资料，联系可可

← 上一章 📖 返回目录 下一章 →