3. 内存带宽瓶颈分析
做量化计算，说白了就是在跟内存带宽较劲。
我刚开始接触这块时，总觉得GPU有几千个核心，算力肯定不是问题。结果一跑量化模型，发现计算单元大部分时间都在“等数据”。嗯，这就是我们今天要聊的核心——内存带宽瓶颈。
3.1 为什么内存带宽是量化计算的首要瓶颈
量化计算有个特点：计算密度低。什么意思？
你想想看，一个INT8的乘加操作，只需要几个时钟周期。但为了拿到这个操作数，你得从显存里搬数据。一次显存访问，可能要几百个时钟周期。这差距，差不多两个数量级。
核心矛盾：计算速度远快于数据搬运速度。
我在项目中遇到过这样一个场景：用FP16跑一个简单的矩阵乘法，显存带宽利用率只有30%左右。换成INT8量化后，按理说计算量减半，应该更快才对。结果呢？性能反而下降了。为什么？因为数据搬运量没变，但计算单元太快了，等数据的时间占比更大。
这就是量化计算的典型困境——计算不再是瓶颈，带宽才是。
3.2 带宽利用率计算
怎么判断你的程序是不是被带宽卡住了？我习惯用两个指标：
- 理论带宽：显卡规格上写的，比如RTX 3090是936 GB/s
- 实际带宽：通过profiler测出来的，比如Nsight Compute
带宽利用率 = 实际带宽 / 理论带宽 × 100%
举个例子：
// 假设你有一个kernel，搬运了1GB数据，耗时2ms
// 实际带宽 = 1GB / 2ms = 500 GB/s
// 如果理论带宽是936 GB/s
// 利用率 = 500 / 936 ≈ 53.4%我的经验：利用率超过80%就算不错了。低于50%的话，你得好好看看访存模式。
我曾经遇到一个项目，带宽利用率只有12%。查了半天，发现是数据在全局内存里是随机访问的，完全没有合并。改成连续访问后，利用率直接跳到75%。
3.3 影响带宽利用率的关键因素
说白了，就三个东西：
- 访存模式——是不是连续的？
- 数据复用——能不能用共享内存缓存？
- 计算密度——每个数据被用了多少次？
3.3.1 访存模式
GPU的显存控制器喜欢连续访问。你想想看，一次内存事务是128字节。如果你只用了其中4个字节，剩下的124字节就浪费了。这就是带宽利用率低的主要原因。
我建议：尽量让每个线程访问连续地址。比如：
// 好的模式：连续访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];
// 坏的模式：跨步访问
int idx = threadIdx.x * gridDim.x + blockIdx.x;
float val = data[idx]; // 每次访问都跨一大段3.3.2 数据复用
量化计算里，权重和激活值经常被重复使用。比如卷积操作，同一个权重会被多个输出像素用到。
这时候，共享内存就派上用场了。把数据先搬到共享内存里，再从那里读，速度能快一个数量级。
注意：共享内存只有几十KB，别塞太多。我见过有人把整个权重矩阵都塞进去，结果bank conflict严重，反而更慢。
3.3.3 计算密度
计算密度 = 计算量 / 数据搬运量。单位是FLOP/Byte。
量化计算的计算密度通常很低。比如一个INT8矩阵乘法，每个元素只做一次乘加，但要从显存里读两次（权重和输入）。计算密度大概只有0.5 FLOP/Byte左右。
而GPU的“甜点”计算密度，一般在10 FLOP/Byte以上。差距很大。
3.4 优化策略
针对带宽瓶颈，我常用的优化手段有这几个：
| 优化手段 | 适用场景 | 效果 |
|---|---|---|
| 数据打包 | INT4/INT8量化 | 减少数据搬运量 |
| 共享内存tiling | 卷积、矩阵乘法 | 提高数据复用 |
| 向量化访存 | 连续数据访问 | 提高带宽利用率 |
| 异步拷贝 | 数据预取 | 隐藏访存延迟 |
3.4.1 数据打包
量化后数据变小了，但显存事务还是128字节。所以，尽量把多个量化数据打包到一个32位或64位变量里。
// 4个INT8打包到一个int32里
int32_t packed = (int32_t)val0 | ((int32_t)val1 << 8)
| ((int32_t)val2 << 16) | ((int32_t)val3 << 24);
// 读取时解包
int32_t data = input[idx];
int8_t v0 = data & 0xFF;
int8_t v1 = (data >> 8) & 0xFF;
// ...避坑指南：我曾经在打包时没注意字节序，结果在x86和ARM上跑出不同结果。嗯，后来统一用little-endian才搞定。
3.4.2 共享内存tiling
这是最经典的优化手段。把数据分块，每块加载到共享内存里，然后所有线程共享使用。
我习惯把tile大小设为32×32或16×16。太小了复用率低，太大了共享内存不够用。
3.4.3 向量化访存
用float4或int4来读取数据。一次读16字节，比读4次4字节效率高得多。
// 向量化读取
float4 vals = reinterpret_cast<float4*>(input)[idx];
// 等价于4次标量读取，但带宽利用率更高3.5 带宽瓶颈分析流程图
下面这张图，是我自己总结的带宽瓶颈分析流程。每次遇到性能问题，我都会走一遍这个流程。
3.6 实战建议
最后，给你几个我踩过坑后的建议：
- 先测带宽利用率，别上来就优化。我见过有人花了两周优化计算逻辑，结果瓶颈在带宽上，白干了。
- 量化后一定要重新测。量化改变了计算密度，带宽瓶颈会更明显。
- 别迷信理论带宽。实际能跑到70%就算优秀了。我曾经在某个老架构上，死活只能跑到40%。
警告：不要为了优化带宽而牺牲正确性。我见过有人为了对齐内存，把数据边界处理错了，结果精度掉了两个数量级。
嗯，带宽瓶颈这块，说白了就是“数据搬得慢”。理解了这一点，量化计算的优化方向就清晰了——要么少搬数据，要么搬得更快。