GPU、Roofline 与异构瓶颈 / GPU, Roofline and Heterogeneous Bottlenecks
GPU 性能分析最容易走偏。
很多人一看到 CUDA Kernel，就直接打开 Nsight Compute 研究 Occupancy、寄存器、Shared Memory 和 Warp Stall。可如果端到端时间主要花在 CPU 准备、数据传输或同步等待，单个 Kernel 再优化也救不了整条流水线。
所以 GPU 性能要分两层看：
端到端异构流水线
-> CPU、GPU、拷贝、同步、队列是否顺畅
单个 GPU Kernel
-> 访存、计算、并发、指令和资源是否高效1. 先看流水线，再看 Kernel
一个 GPU 项目通常不是只有 Kernel。
CPU 准备输入
|
v
Host -> Device 拷贝
|
v
Kernel A
|
v
Kernel B
|
v
Device -> Host 拷贝
|
v
CPU 后处理慢可能发生在任何位置。
Nsight Systems 用来看整条时间线。
Nsight Compute 用来看单个 Kernel。
顺序通常是：
Nsight Systems
|
+-- GPU 是否空闲？
+-- 拷贝和计算是否重叠？
+-- Kernel 之间是否有空洞？
+-- CPU 是否在等待 GPU？
+-- GPU 是否在等待 CPU 提交？
|
v
选出值得深入的 Kernel
|
v
Nsight Compute
|
+-- 内存访问
+-- Occupancy
+-- Warp Stall
+-- 吞吐上限如果 Systems 显示 GPU 大量空闲，Compute 的局部优化优先级就很低。
2. Roofline 的直觉
Roofline 是一个性能上限模型。
它把程序放在“计算能力”和“内存带宽”两个屋顶下面。
核心概念是 Arithmetic Intensity，算术强度。
Arithmetic Intensity = 计算量 / 数据搬运量例如每读写很多字节只做一次加法，算术强度低。
例如矩阵乘法中每个数据被重复参与大量乘加，算术强度高。
简化图：
Performance
^
| compute roof
|-------------------------------
| /
| /
| /
| /
| memory roof /
+----------------------------------> Arithmetic Intensity左侧受内存带宽限制。
右侧受计算峰值限制。
Roofline 的意义不是精确预测每毫秒，而是阻止错误努力。
如果一个 Kernel 明显受内存带宽限制，你继续手写更多计算指令通常没用。应该先减少数据搬运、改善合并访问、复用数据或改变算法。
3. GPU 时间线常见瓶颈
GPU 空洞
GPU 时间线上有大段空白，说明 GPU 没有工作。
原因可能是：
- CPU 准备数据太慢；
- Kernel 启动太频繁且粒度太小；
- 同步点过多；
- Stream 没有并发；
- 数据传输阻塞；
- 依赖关系串行；
- 内存分配发生在热路径；
- 任务调度器没有及时提交。
拷贝和计算没有重叠
Host 和 Device 之间传输数据可能占很大时间。
如果每次都同步拷贝再执行 Kernel，再同步拷贝回来，GPU 会经常等待。
优化方向包括 pinned memory、异步拷贝、多 Stream、批处理和减少 Host/Device 往返。
频繁同步
同步会切断流水线。
常见同步来源：
cudaDeviceSynchronize()；- 同步内存拷贝；
- 默认 Stream 的隐式依赖；
- 读取 Device 结果到 Host；
- 错误检查写法过于粗；
- 调试或日志路径。
同步不是不能用，而是要知道它放在哪里、阻断了什么。
4. Nsight Systems 的使用方式
Nsight Systems 的目标是建立端到端图景。
你可以先加 NVTX 标记业务阶段。
nvtxRangePushA("build input");
buildInput();
nvtxRangePop();
nvtxRangePushA("solve kernels");
launchKernels();
nvtxRangePop();这样时间线里不只有 CUDA API 和线程号，还能看到项目语义。
常见观察顺序：
1. 先看总时间和业务阶段
2. 看 CPU 线程是否有长等待
3. 看 CUDA API 调用是否密集
4. 看 GPU 队列是否有空洞
5. 看 memcpy 和 kernel 是否重叠
6. 看同步点是否切断并发
7. 选出一个真正影响端到端的 kernel命令形态可以是：
nsys profile -t cuda,nvtx,osrt -o report ./app
nsys stats report.nsys-repcuda 记录 CUDA API 和 GPU 工作。
nvtx 记录应用标记。
osrt 记录部分操作系统运行库事件，例如锁和 I/O 调用。
5. Nsight Compute 的使用方式
Nsight Compute 针对单个 Kernel。
它回答：
- 一个 Kernel 的 DRAM/L2/Shared Memory 行为如何；
- Warp 为什么停顿；
- Occupancy 受寄存器、Shared Memory 还是线程块限制；
- 指令吞吐是否接近硬件上限；
- 内存访问是否合并；
- 分支是否导致发散；
- 每个 Source Line 或 SASS 指令的成本在哪里。
命令形态：
ncu --set full ./app
ncu --kernel-name my_kernel ./app不要对整个复杂程序无脑 --set full。
它可能很慢，数据也很多。通常先用 Nsight Systems 找到目标 Kernel，再让 Nsight Compute 只采这个 Kernel。
6. Occupancy 不是最终目标
Occupancy 描述 GPU 上可驻留 Warp 的比例。
它低时可能表示寄存器、Shared Memory 或线程块配置限制并发。
但 Occupancy 高不等于性能好。
如果 Kernel 受内存带宽限制，高 Occupancy 可能只是让更多 Warp 一起等内存。
如果 Kernel 指令依赖很强，增加 Occupancy 可能能隐藏延迟。
如果 Kernel 每个线程寄存器很多，强行降低寄存器使用可能导致 spilling，把数据溢出到本地内存，反而变慢。
所以 Occupancy 只能作为解释指标，不能作为目标指标。
7. GPU 访存问题
GPU 访存常见问题：
- 全局内存访问不合并；
- 访问 stride 太大；
- 读写重复；
- Shared Memory bank conflict；
- 寄存器 spilling；
- Host/Device 频繁往返；
- Unified Memory 迁移页；
- 数据布局不适合线程组织；
- 原子操作竞争。
常见优化方向：
- 改善线程到数据的映射；
- 使用 SoA 而不是 AoS；
- 分块和 Shared Memory 复用；
- 减少全局内存写；
- 合并多个小 Kernel；
- 使用异步拷贝和流水线；
- 降低原子热点；
- 把 CPU 后处理移到 GPU 或反过来减少往返。
8. CPU/GPU 共同瓶颈
异构程序常见问题不是某一个设备慢，而是协作方式慢。
CPU: [prepare][wait gpu....][prepare][wait gpu....]
GPU: [idle...][kernel][idle...][kernel]这说明 CPU 和 GPU 没有形成流水线。
可能的优化方向：
- 把任务批量化；
- 减少同步；
- 使用多个 Stream；
- 用双缓冲隐藏拷贝；
- 让 CPU 准备下一批数据；
- 延迟读取结果；
- 把小 Kernel 合并；
- 把部分轻量任务留在 CPU，避免传输开销。
但要用端到端指标判断。GPU 利用率提高，如果总延迟变差，仍然不是成功。
9. Roofline 报告模板
目标：
pressure_update kernel 占端到端 31%，需要判断是否值得优化。
Systems 证据：
GPU 在求解阶段利用率高，CPU 准备不是主瓶颈。
Compute 证据：
DRAM 吞吐接近硬件可达带宽；
每个元素执行少量计算；
L2 命中率低，访问模式为 stride。
模型：
算术强度低，属于 memory-bound。
实验：
调整数据布局，让相邻线程读取连续地址。
预期：
Global load efficiency 上升，DRAM 字节数下降或有效吞吐上升。
成功标准：
Kernel 时间下降，并且端到端求解时间下降。10. 本篇总结
GPU 性能要先看异构流水线，再看单个 Kernel。
Nsight Systems 帮你判断 CPU、GPU、拷贝和同步之间的时间关系。
Nsight Compute 帮你解释某个 Kernel 内部的访存、计算、Occupancy 和停顿。
Roofline 模型帮助你判断应该减少数据搬运、提高复用、改善访问模式，还是追求更高计算吞吐。