6.9 面试追问
KV 块分配、分页寻址与前缀复用
「KV Cache 量化后为什么端到端反而变慢?」
因为反量化路径不统一。好的做法是 Decode 融合反量化(在 Paged Attention Kernel 里读 INT8 → 转 FP32 → 乘 scale → 直接进 online softmax,不物化浮点 Cache);但如果 Prefill 走的是「物化 gather + 反量化」(按 Block Table 把 KV 还原成浮点张量再交给 SDPA(Scaled Dot-Product Attention,PyTorch 官方提供的注意力算子封装)),每层都要重新 gather + 反量化 + 写一遍浮点张量,这个开销会吃掉带宽收益。结论:KV 量化的收益取决于反量化是否融合。
码 6-2 · 反量化的两条路径:融合 vs 物化
# 路径 A(好):融合在 Paged Attention kernel 内部,不物化浮点 Cache
key = tl.load(k_cache_ptr + k_offsets, mask=mask, other=0.0
).to(tl.float32) * k_scale[:, None]
# → 直接进 online softmax,全程只多了一次乘法和一次类型转换
# 路径 B(有代价):按 Block Table 把 KV 物化成浮点张量,再交给 SDPA
values = cache[block_ids, block_offsets].to(torch.float32) # [seqlen, H_kv, D]
token_scales = scales[block_ids, block_offsets].float()
k_fp = (values * token_scales.unsqueeze(-1)).to(dtype)
# → 每一层都要重新 gather + 反量化 + 写一遍浮点张量
同样是把 INT8 还原成浮点,两条路径的成本差一个量级:路径 A 只在寄存器里过一遍,带宽收益完整保留;路径 B 要额外分配、写入、再读出
[seqlen, H_kv, D]的浮点张量,把省下的带宽又还回去了。所以「KV 量化端到端到底快不快」,答案取决于 Prefill 侧走的是哪条路——这也是「量化收益取决于反量化是否融合」这句话的代码级解释。
「为什么带宽减半,Kernel 只快一点?」
反量化把 K/V tile 提到 FP32 精度,寄存器需求上升(往往需要更多 warp),ALU(Arithmetic Logic Unit,算术逻辑单元) 与寄存器压力抵消了一部分带宽收益。进一步优化方向是代数等价改写:先算
k_int8 · q的点积,再乘 scale 向量,而不是先缩放整个 tile——乘法次数可从 O(block × head_dim) 降到 O(block)。