3. CUDA kernel性能指标：occupancy、throughput、latency、memory bandwidth、compute utilization

在CUDA kernel的性能调优中，我们经常听到“占用率不够高”“带宽没打满”“延迟掩盖不住”等说法。这些判断背后，其实是一组相互关联、又各有侧重的性能指标。理解这些指标的含义、测量方法以及它们之间的权衡关系，是定位GPU kernel瓶颈的第一步。本章将围绕occupancy（占用率）、throughput（吞吐量）、latency（延迟）、memory bandwidth（内存带宽）和compute utilization（计算利用率）五个核心指标展开，帮助你建立一套系统化的性能分析框架。

上图展示了五个核心指标之间的逻辑关系。在实际调优中，我们通常先观察计算利用率和内存带宽的数值，判断kernel是计算密集型还是访存密集型，再结合占用率和延迟数据，确定具体的优化方向。下面逐一展开每个指标的原理、测量方法和实战意义。

3.1 Occupancy（占用率）

核心原理

Occupancy定义为每个流多处理器（SM）上活跃的warp数量与硬件支持的最大warp数量之比。它反映了GPU隐藏延迟的能力——当某些warp因访存或同步而等待时，调度器可以立即切换到其他就绪的warp继续执行。高占用率意味着有更多的warp可供调度，从而更有效地掩盖内存延迟和指令延迟。

占用率受多个因素限制：寄存器数量、共享内存大小、线程块维度、硬件资源上限（如SM最大线程数、最大warp数）。例如，一个SM最多支持64个warp（2048线程），如果每个线程使用32个寄存器，而SM寄存器文件为65536个，则最多容纳2048个线程，占用率可达100%；但如果每个线程使用40个寄存器，则最多容纳1638个线程，占用率降至80%。

判断方法

使用NVIDIA Nsight Compute（ncu）或CUDA Occupancy Calculator可以快速获取占用率数据。在ncu中，执行以下命令：

ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./your_kernel

该指标直接给出kernel执行期间的平均活跃warp占比。此外，还可以查看launch__occupancy_limit_registers、launch__occupancy_limit_shared_mem等指标，了解是哪个资源限制了占用率。

案例拆解

假设一个kernel的占用率仅为40%，通过ncu发现限制因素是寄存器使用量（每个线程48个寄存器）。优化方案：使用__launch_bounds__(256, 8)限制每个线程块256线程、每SM至少8个块，编译器会强制减少寄存器分配，可能引入少量局部内存溢出（spill），但整体占用率提升到75%，kernel执行时间缩短了22%。

3.2 Throughput（吞吐量）

核心原理

Throughput指单位时间内kernel完成的工作量，通常以每秒处理的元素数、每秒执行的指令数（IPS）或每秒浮点操作数（FLOPS）来衡量。它反映的是GPU整体处理能力是否被充分利用。与延迟不同，吞吐量关注的是“流水线”的流量，而非单个操作的耗时。

在GPU中，吞吐量受限于三个主要瓶颈：计算单元（CUDA Core/Tensor Core）的峰值吞吐、内存子系统的带宽、以及指令发射带宽。一个kernel的实际吞吐量由最弱的环节决定。

判断方法

使用ncu获取以下指标：

ncu --metrics sm__inst_executed.avg.per_cycle_active,sm__throughput.avg.pct_of_peak_sustained_elapsed ./your_kernel

其中sm__throughput.avg.pct_of_peak_sustained_elapsed表示SM整体吞吐量占峰值的百分比。如果该值接近100%，说明SM计算单元已饱和；如果远低于100%，则需要检查是否存在访存瓶颈或调度开销。

常见误区

一个常见误区是认为“吞吐量低就一定是计算单元不够快”。实际上，很多kernel的吞吐量受限于内存带宽或指令发射带宽。例如，一个简单的向量加法kernel，其计算吞吐量可能只有峰值的10%，但内存带宽利用率已达90%，此时优化方向应转向减少内存访问量（如使用向量化加载、共享内存复用），而非增加计算并行度。

3.3 Latency（延迟）

核心原理

Latency指从发出一个操作（如加载数据、执行指令）到该操作完成所需的时间。在GPU中，主要关注两类延迟：内存延迟（global memory访问延迟，通常400-800个周期）和指令延迟（算术指令延迟，通常4-20个周期）。GPU通过大量并发warp来掩盖延迟——当一个warp等待内存返回时，调度器切换到其他warp执行。

延迟掩盖能力取决于两个因素：占用率（活跃warp数量）和每个warp的延迟容忍度（即warp在等待时是否还有其他可执行的指令）。如果kernel中每个warp只有一条加载指令后紧跟着一条使用该数据的指令，那么即使占用率很高，也可能出现大量等待周期。

判断方法

使用ncu的smsp__average_warps_issue_stalled_*系列指标，可以查看warp停滞的原因分布：

ncu --metrics smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio,smsp__average_warps_issue_stalled_short_scoreboard_per_issue_active.ratio ./your_kernel

其中long_scoreboard表示等待全局内存返回，short_scoreboard表示等待共享内存或常量内存。如果long_scoreboard占比很高，说明内存延迟掩盖不足。

实战场景

考虑一个矩阵转置kernel，每个线程读取一行数据并写入转置位置。由于写入地址不连续，导致内存访问模式为“读连续、写离散”，产生大量long_scoreboard停滞。优化方案：使用共享内存作为中转，先按行读取到共享内存，再按列写入全局内存。修改后，long_scoreboard停滞比例从65%降至20%，kernel耗时减少40%。

3.4 Memory Bandwidth（内存带宽）

核心原理

内存带宽指GPU全局内存在单位时间内可传输的数据量，单位通常为GB/s。现代GPU（如A100）的峰值带宽可达2TB/s以上。实际带宽利用率是衡量访存密集型kernel效率的关键指标。

影响带宽利用率的因素包括：内存访问模式（是否合并访问）、数据宽度（128位向量化加载）、缓存命中率、以及访问地址的对齐情况。理想情况下，同一warp的32个线程应访问连续的128字节内存段，实现一次事务完成所有数据传输。

判断方法

使用ncu获取以下指标：

ncu --metrics dram__bytes_read.sum.per_second,dram__bytes_write.sum.per_second,gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed ./your_kernel

gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed直接给出内存吞吐量占峰值的百分比。如果该值超过80%，说明内存子系统接近饱和，kernel为访存密集型。

案例拆解

一个图像滤波kernel，每个线程处理一个像素，读取3x3邻域数据。初始实现中，每个线程独立读取9个像素，导致大量非合并访问，带宽利用率仅35%。优化方案：使用共享内存加载图像块（tile），每个线程只从全局内存读取一次数据，后续邻域访问从共享内存获取。优化后带宽利用率提升至82%，kernel耗时减少55%。

3.5 Compute Utilization（计算利用率）

核心原理

计算利用率指SM中计算单元（CUDA Core、Tensor Core、SFU等）在kernel执行期间的平均活跃比例。它反映了计算资源是否被充分利用。与吞吐量不同，计算利用率更关注“计算单元的空闲程度”，而非整体工作量。

计算利用率低通常意味着kernel存在访存等待、同步开销、或指令级并行度不足。例如，一个kernel中浮点运算之间穿插了大量整数运算和分支跳转，导致CUDA Core的浮点单元利用率下降。

判断方法

使用ncu查看不同计算单元的利用率：

ncu --metrics sm__inst_executed_pipe_fma.avg.pct_of_peak_sustained_elapsed,sm__inst_executed_pipe_alu.avg.pct_of_peak_sustained_elapsed ./your_kernel

如果FMA（浮点乘加）管道利用率高而ALU（整数/逻辑）管道利用率低，说明kernel以浮点计算为主；反之则可能包含大量地址计算或循环控制逻辑。

常见误区与总结

一个常见误区是“计算利用率越高越好”。实际上，对于访存密集型kernel，计算利用率低是正常现象，因为瓶颈在内存而非计算。此时强行提高计算利用率（如增加冗余计算）反而会降低性能。正确的做法是：先判断kernel类型（计算密集型/访存密集型/混合型），再针对瓶颈进行优化。