计算瓶颈分析：识别计算密集型任务、内存带宽瓶颈、延迟隐藏技术

好，咱们直接进入正题。做核函数优化，第一步不是动手改代码，而是搞清楚——你的程序到底慢在哪？

我见过太多人，一上来就调循环、改指令，结果折腾半天性能没提升。为什么？因为没找准瓶颈。这就像看病，你得先诊断，再开药。

这一章，我们就来聊聊怎么诊断。说白了，就是三件事：找计算密集的地方、看内存带宽够不够、想办法隐藏延迟。

1. 计算密集型任务：你的GPU在“算”还是在“等”？

先问个问题：你的核函数里，算术操作多，还是内存访问多？

如果算术操作占大头，那就是计算密集型。比如矩阵乘法、FFT、卷积——这些任务里，GPU的ALU（算术逻辑单元）忙得飞起，但内存访问相对少。

我个人习惯，判断一个核函数是不是计算密集型，会先看算术强度。公式很简单：

算术强度 = 总浮点操作数 / 总字节访问量

单位是 FLOP/Byte。这个值越高，说明计算越密集。

举个例子。我做过一个图像滤波的核函数，每个像素要做几十次乘加，但只读几个像素值。算术强度轻松超过 10 FLOP/Byte。这种任务，GPU 的算力就是瓶颈。

嗯，这里要注意：别只看理论峰值。实际运行时，指令发射、寄存器压力、分支分歧都会影响。我建议你用 nvprof 或 Nsight Compute 跑一下，看 sm_efficiency 和 achieved_occupancy。如果 sm_efficiency 接近 100%，但 occupancy 很低，那大概率是计算瓶颈。

2. 内存带宽瓶颈：数据搬运才是“真·慢”

说实话，大部分核函数其实不是算得慢，而是数据搬得慢。

你想想看，GPU 的算力动辄几十 TFLOPS，但显存带宽才几百 GB/s。一个简单的加法：读两个数、加一下、写回去。算术操作只有 1 FLOP，但内存访问是 12 字节（两个 4 字节读 + 一个 4 字节写）。算术强度只有 0.083 FLOP/Byte——典型的带宽瓶颈。

我在项目中遇到过最典型的例子：一个粒子模拟的核函数，每个粒子只做一次位置更新（几行代码），但要从全局内存读位置、速度，再写回去。结果呢？GPU 的算力利用率不到 5%，大部分时间都在等数据。

怎么识别？看两个指标：

指标	正常范围	带宽瓶颈信号
DRAM 利用率	60-80%	> 90% 且计算单元空闲
L2 命中率	40-60%	< 20% 说明访存模式差
算术强度	> 1 FLOP/Byte	< 0.5 FLOP/Byte

解决带宽瓶颈，核心思路就两个：减少访问量（比如用共享内存缓存复用数据）和合并访问（让相邻线程访问连续地址）。

3. 延迟隐藏技术：让GPU“忙”起来

好，现在你知道瓶颈在哪了。但还有一个问题：即使你解决了计算和带宽，GPU 还是可能“偷懒”。

为什么？因为延迟。一条指令从发出到执行完，需要几十个时钟周期。如果 GPU 在这期间没事干，那就浪费了。

延迟隐藏，说白了就是让 GPU 在等数据的时候，去干别的活。怎么干？靠线程切换。

GPU 的 SM（流多处理器）里有很多 warp（线程束）。当一个 warp 在等内存数据时，SM 会立刻切换到另一个 warp 执行。只要 warp 够多，延迟就能被“藏”起来。

我个人的经验是：occupancy（占用率）越高，延迟隐藏越好。但也不是越高越好。我记得有一次，我把 occupancy 从 50% 提到 100%，性能反而下降了。为什么？因为每个线程分到的寄存器太少，导致寄存器溢出到本地内存，反而更慢。

还有一个技巧：指令级并行。在同一个线程里，把独立的计算和内存访问交错开。比如：

// 不好的写法：先等数据，再计算
float a = data[i];
float b = data[i+1];
float c = a + b;  // 这里在等 a 和 b

// 好的写法：提前发出内存请求
float a = data[i];
float b = data[i+1];
float c = a + b;  // 编译器可能自动优化，但手动写更可靠
// 或者用 __ldg() 强制使用只读缓存

嗯，这里要提醒一下：别过度优化。现代编译器（比如 nvcc）已经能做很多指令重排。你手动写一堆复杂的交错代码，可能反而让编译器优化不了。

知识体系总览

最后，我用一张图总结一下这一章的核心逻辑。你一看就明白了：

这张图把三个核心概念串起来了。你从左往右看：先判断是不是计算密集型，如果不是，再看内存带宽，最后用延迟隐藏技术兜底。每一步都有对应的指标和优化方向。

好了，这一章就到这里。记住：别急着改代码，先拿数据说话。下一章我们会深入具体的优化技巧，到时候再细聊。

---