3、nvprof高级功能：自定义范围标记（nvtx）、指标收集（--metrics）、追踪API调用（--trace）

好，咱们继续深入。上一章我们把nvprof的基础用法捋了一遍，说白了就是能看个大概。但真正做性能调优，光看个大概可不够。你得能精准定位到代码里的某个函数、某段循环，甚至某一行计算。这就是我今天要聊的三个高级功能：NVTX标记、指标收集、API追踪。

我个人习惯，拿到一个CUDA程序，第一件事不是跑profiling，而是先想清楚：我想看什么？ 是核函数的执行时间？是内存拷贝的带宽？还是API调用的开销？想清楚目标，再选工具。这三个功能，正好对应了三种不同的分析视角。

3.1 自定义范围标记（NVTX）

先说说NVTX。这玩意儿全称是NVIDIA Tools Extension SDK。名字挺唬人，其实说白了就是给你在代码里插旗子用的。

你想想看，一个复杂的CUDA程序，可能有几十个核函数，几十次内存拷贝。nvprof默认会把所有事件都列出来，但你根本分不清哪个是哪个。这时候，NVTX就派上用场了。

怎么用？ 很简单，三步走：

包含头文件 nvtx3/nvToolsExt.h
在你想标记的代码段前后，加上 nvtxRangePush 和 nvtxRangePop
运行nvprof时，加上 --nvtx 选项

来看个例子：

#include <nvtx3/nvToolsExt.h>

void myKernel() {
    // 标记开始
    nvtxRangePushA("数据预处理阶段");
    
    // 你的CUDA代码...
    kernel1<<<grid, block>>>(d_data);
    cudaMemcpy(d_result, h_result, size, cudaMemcpyDeviceToHost);
    
    // 标记结束
    nvtxRangePop();
    
    // 另一个阶段
    nvtxRangePushA("核心计算阶段");
    kernel2<<<grid, block>>>(d_result);
    nvtxRangePop();
}

然后运行：

nvprof --nvtx ./myapp

输出结果里，所有事件就会按照你标记的范围分组。一眼就能看出哪个阶段耗时最多。

我的小技巧： 标记名称别用太长的字符串，尽量用英文缩写。比如"preproc"代替"数据预处理阶段"。为什么？因为nvprof的GUI界面显示宽度有限，太长的名字会被截断，反而看不清。

我在项目中遇到过一个问题：有个同事在循环里反复调用 nvtxRangePush 和 nvtxRangePop，结果导致性能下降。嗯，这里要注意——NVTX本身是有开销的，虽然很小，但如果你在热路径（比如循环体内部）频繁调用，还是会有影响。我建议只在函数入口和出口、或者关键阶段切换处使用。

3.2 指标收集（--metrics）

接下来是指标收集。这个功能，说白了就是让你能看硬件的底层数据。

默认情况下，nvprof只给你看时间。但时间只是表象。为什么慢？是计算单元没跑满？还是内存带宽不够？还是指令流水线有停顿？这些问题的答案，藏在硬件计数器里。

用法：

nvprof --metrics all ./myapp   # 收集所有指标（不推荐，太慢了）
nvprof --metrics gld_efficiency,gst_efficiency ./myapp  # 只收集你关心的

常用的指标有哪些？我列个表：

指标名称	含义	什么时候看
gld_efficiency	全局内存加载效率	怀疑内存访问模式有问题时
gst_efficiency	全局内存存储效率	同上
occupancy	SM占用率	怀疑计算单元没跑满时
sm_efficiency	SM活跃度	看整体利用率
dram_read_throughput	显存读取带宽	看是否达到硬件上限
inst_per_warp	每个warp执行的指令数	看代码是否有冗余指令

核心思路： 先看 sm_efficiency。如果这个值很低（比如低于50%），说明GPU大部分时间在空转。这时候再看 occupancy 和 gld_efficiency，就能定位到是线程块太小，还是内存访问没对齐。

我曾经调过一个图像处理的程序，sm_efficiency只有30%。我一开始以为是计算量不够，结果一看gld_efficiency，只有15%。嗯，问题出在内存访问上。后来把数据布局从AoS改成SoA，gld_efficiency直接飙到85%，sm_efficiency也到了70%。

注意： 收集指标会显著增加运行时间。因为nvprof需要多次运行程序来采集不同的计数器。我建议你先用 --metrics 指定少量指标，定位问题后，再用更细粒度的指标深入分析。别一上来就 --metrics all，那会等到你怀疑人生。

3.3 追踪API调用（--trace）

最后一个功能，API追踪。这个功能，说白了就是看你的程序调用了哪些CUDA API，以及每次调用的耗时。

为什么要看这个？因为很多时候，性能瓶颈不在核函数里，而在API调用上。比如频繁的 cudaMalloc 和 cudaFree，或者没必要的 cudaMemcpy。

用法：

nvprof --trace api ./myapp   # 追踪所有API调用
nvprof --trace cudaMemcpy ./myapp  # 只追踪内存拷贝

输出结果会列出每个API调用的名称、耗时、调用次数。我一般会关注两个点：

调用次数： 如果某个API被调用了成千上万次，那肯定有问题。比如 cudaMalloc 应该在初始化时一次性分配，而不是在循环里反复分配。
单次耗时： 比如 cudaMemcpy 如果每次只拷贝几个字节，那开销主要来自API调用本身，而不是数据传输。

来看个实际输出：

Time(%)  Time     Calls   Avg     Min     Max  Name
 45.2%  2.345s    1000  2.345ms 1.2ms  5.6ms  cudaMemcpy
 30.1%  1.567s     500  3.134ms 2.1ms  4.8ms  cudaMalloc
 24.7%  1.283s     200  6.415ms 3.5ms  9.2ms  cudaLaunchKernel

看到没？cudaMemcpy 被调用了1000次，占了将近一半的时间。这时候你就该想想：能不能合并内存拷贝？或者用异步拷贝 + 流水线来隐藏延迟？

我的经验： 我建议把 --trace api 和 --trace gpu 结合起来用。前者看API调用的开销，后者看GPU执行的开销。有时候你会发现，API调用花了很长时间，但GPU实际执行时间很短。这说明你的程序被CPU端的API调用拖慢了。

我曾经遇到过一个案例：一个深度学习训练程序，每次迭代都要调用 cudaSetDevice 来切换GPU。结果这个API本身耗时只有几微秒，但因为它会触发上下文切换，导致后续的核函数启动延迟暴增。后来我们把设备绑定放到初始化阶段，性能直接提升了15%。

知识体系总览

这三个功能，其实对应了性能分析的三个层次：

这三个功能，我建议你按这个顺序用：

先用 --trace api 看看有没有明显的API滥用
再用NVTX标记关键代码段，缩小分析范围
最后用 --metrics 深入分析硬件层面的瓶颈

这样一层层剥下去，再复杂的性能问题也能找到根因。

避坑指南： 我曾经在调试一个多GPU程序时，同时开启了 --metrics 和 --trace，结果nvprof跑了将近半小时还没出结果。后来才发现，这两个选项会互相影响，导致nvprof需要多次重跑程序。我的建议是：分开使用，先trace后metrics，别贪心。

好了，这一章的内容就到这里。这三个功能，你只要用熟了，基本上就能应付90%的性能分析场景。记住，工具是死的，思路是活的。别光顾着收集数据，要多想想数据背后的含义。

📚 课程目录

3、nvprof高级功能：自定义范围标记（nvtx）、指标收集（--metrics）、追踪API调用（--trace）

3.1 自定义范围标记（NVTX）

3.2 指标收集（--metrics）

3.3 追踪API调用（--trace）

知识体系总览