4. CUDA性能瓶颈定位：计算瓶颈、访存瓶颈、延迟隐藏、带宽利用率分析

性能调优这件事，说白了就是找短板。GPU 计算也一样——你得先知道瓶颈在哪，才能对症下药。我见过不少同学一上来就调 block 大小、改 shared memory，折腾半天发现根本没打在点上。嗯，今天我们就来系统地聊聊，怎么定位 CUDA 程序的性能瓶颈。

4.1 计算瓶颈 vs 访存瓶颈：先搞清楚谁在拖后腿

一个 kernel 跑得慢，无非两个原因：要么是计算单元在等数据（访存瓶颈），要么是数据在等计算单元（计算瓶颈）。怎么区分？看 occupancy 和指令吞吐就够了。

我个人习惯先看 算术强度（Arithmetic Intensity）。这个指标很简单：

算术强度 = 总浮点运算次数 / 总访存字节数

如果算术强度很低（比如 < 1 FLOP/byte），那大概率是访存瓶颈。反之，如果算术强度很高（比如 > 10 FLOP/byte），那计算瓶颈的可能性更大。

举个例子，我在优化一个矩阵乘法时，发现算术强度只有 0.8。当时我就知道，这 kernel 肯定被内存带宽卡死了。后来改用 tiling 和 shared memory，算术强度提升到 12，性能直接翻了三倍。

4.2 延迟隐藏：GPU 的“障眼法”

GPU 为什么能容忍高延迟？因为它靠的是 大量线程切换 来隐藏延迟。你想想看，当一个 warp 在等数据时，SM 会立刻切换到另一个 warp 执行。只要 warp 够多，延迟就被“藏”起来了。

但这里有个坑——occupancy 不是越高越好。我曾经在一个项目中，把 occupancy 从 50% 推到 100%，结果性能反而下降了。为什么？因为每个线程分到的寄存器变少了，编译器不得不 spill 到 local memory，反而更慢。

延迟隐藏的效果，可以用 理论占用率 和 实际占用率 的差距来衡量。差距越小，说明隐藏得越好。

4.3 带宽利用率：别让显存闲着

访存瓶颈的核心指标就是 带宽利用率。说白了，就是你的程序实际用到了多少显存带宽。

用 Nsight Compute 看 l1tex__throughput 和 dram__throughput，如果利用率低于 60%，那就有优化空间。

我记得有一次调一个图像处理 kernel，带宽利用率只有 25%。一看 profile 结果，发现全局内存访问模式是 非合并访问（non-coalesced）。每个 warp 的 32 个线程访问的地址不连续，导致一次内存事务只用了 1/4 的带宽。

怎么解决？调整数据布局，让相邻线程访问相邻地址。比如把结构体数组（AoS）改成数组结构体（SoA）：

// 不推荐：AoS 布局
struct Particle {
    float x, y, z;
};
Particle particles[N];

// 推荐：SoA 布局
struct Particles {
    float *x, *y, *z;
};
Particles particles;

改完之后，带宽利用率直接飙到 85%。

4.4 实战：用 Nsight Compute 定位瓶颈

工具方面，我强烈推荐 Nsight Compute（ncu）。它比 nvprof 强太多了，能直接告诉你瓶颈类型。

跑一下这个命令：

ncu --set full -o my_kernel_profile ./my_app

然后打开生成的 .ncu-rep 文件，看 Speed of Light 图。它会显示你的 kernel 离理论峰值有多远：

• Memory Throughput：如果接近 100%，说明访存是瓶颈
• Compute Throughput：如果接近 100%，说明计算是瓶颈
• 如果两者都低，说明有 延迟隐藏不足 或 指令发射瓶颈

我曾经优化过一个流体模拟 kernel，Speed of Light 图显示 Memory 只有 30%，Compute 只有 20%。我当时就懵了——两个都低？后来发现是 指令依赖 导致流水线停顿太多。加了 __launch_bounds__ 和手动展开循环后，Compute 利用率升到 70%。

4.5 常见瓶颈速查表

症状	可能原因	排查方法	常用优化
带宽利用率低	非合并访问、bank conflict	检查 l1tex__t_sectors	SoA 布局、对齐访问
occupancy 低	寄存器/共享内存占用过高	查看 sm__occupancy	减少寄存器、拆分 kernel
计算吞吐低	指令依赖、分支发散	检查 sm__inst_executed	循环展开、消除分支
延迟隐藏差	warp 数不足	查看 theoretical occupancy	增加 block 数、减少寄存器

4.6 我的调优流程

最后分享一下我个人的调优套路，你可以参考：

1. 先跑 ncu，看 Speed of Light 图，确定是计算还是访存瓶颈
2. 如果是访存瓶颈，检查内存访问模式、带宽利用率、cache 命中率
3. 如果是计算瓶颈，检查指令吞吐、occupancy、分支发散
4. 如果两者都低，大概率是延迟隐藏问题，增加 warp 数或减少寄存器压力
5. 改完再跑，对比前后指标，确保瓶颈确实被解决了

我曾经犯过一个错误——花了两天优化一个 kernel 的计算部分，结果发现瓶颈其实是访存。嗯，从那以后我再也不凭感觉调优了，一定先看 profile 数据。