执行摘要

• 一句话：修复 BF16 模型使用 turboquant 时 FP8 转换失败
• 推荐动作：此 PR 为典型的 bugfix，变更量小且讨论清晰，值得关注如何通过中间类型转换解决类型兼容性问题。对于维护量化后端的人员有参考价值。

功能与动机

根据 PR body，运行 BF16 模型时 Triton 的 convert_custom_float8 只支持 FP16/FP32 输入，直接 BF16→FP8 触发 AssertionError。修复可让 turboquant 量化正常加载并服务。

实现拆解

1. 在 vllm/v1/attention/ops/triton_turboquant_store.py 的 _tq_fused_store_fp8 函数中，将加载 Key tensor 后的表达式 .to(tl.float32) 添加在 tl.load 之后。
2. 这样无论原始 dtype 是 FP16、BF32 还是 BF16，都会先转成 FP32 再进行 FP8 转换，避免直接 BF16→FP8 的断言错误。
3. 该修改仅增加一个类型转换操作，不影响后续计算路径，且额外开销可忽略。
4. 测试方式：使用 --kv-cache-dtype turboquant_k8v4 启动 Qwen2.5-3B-Instruct，修复前引擎初始化崩溃，修复后服务正常。

关键文件：

• vllm/v1/attention/ops/triton_turboquant_store.py（模块 量化内核；类别 source；类型 bugfix；符号 _tq_fused_store_fp8）: 该文件包含 Turboquant FP8 存储的 Triton 内核，修复后的关键一行代码位于 _tq_fused_store_fp8 函数中。

关键符号：_tq_fused_store_fp8

关键源码片段

vllm/v1/attention/ops/triton_turboquant_store.py

该文件包含 Turboquant FP8 存储的 Triton 内核，修复后的关键一行代码位于 _tq_fused_store_fp8 函数中。

# 在 Triton 内核中加载 Key 缓存并转为 FP8
# 原代码直接对 BF16 值使用 .to(tl.float8e4b15) 导致断言错误
# 修复：先将加载结果转为 FP32（对所有输入类型无损），再转为 FP8
d_offs = tl.arange(0, BLOCK_D)
d_mask = d_offs < D
k_vals = tl.load(Key_ptr + base + d_offs, mask=d_mask, other=0.0).to(tl.float32) # 新增 .to(tl.float32)
k_fp8 = k_vals.to(tl.float8e4b15) if FP8_E4B15 else k_vals.to(tl.float8e4nv)
k_bytes = k_fp8.to(tl.uint8, bitcast=True)
tl.store(KV_cache_ptr + slot_base + d_offs, k_bytes, mask=d_mask)

评论区精华

审查者 vibhavagarwal5 指出此 PR 与 #39908 重复，要求两位作者协调合并。作者 XuZhou26 与 hoseung2 沟通后，一致认为本 PR 方案更简单，决定合并此 PR 并关闭另一条。之后作者多次触发 CI 均因无关失败，最终由 mgoin 批准合并。

• 与重复 PR #39908 的协调合并 (design): 合并当前 PR，关闭另一个。
• 修复方案：无条件转 FP32 是否安全 (correctness): 采用 FP32 中间转换。

风险与影响

• 风险：此变更仅影响 turboquant 注意力算子的 FP8 转换路径，将 BF16 先转为 FP32 再转 FP8，不存在精度或性能退化风险。对所有支持的 dtype 均无损。但若后续有更高 dtype（如 FP64）也不受影响，因为 .to(tl.float32) 会截断，但当前场景无此情况。整体风险极低。
• 影响：影响范围：使用 BF16 模型并启用 turboquant 量化的用户（如 Qwen2 系列）。修复后这些用户可正常使用 turboquant，其他用户无影响。对系统性能无显著改变。
• 风险标记：低风险变更, BF16兼容性

关联脉络

• PR #39908 Fix turboquant BF16 cast (duplicate): 修复相同问题，但方案不同；作者协调后合并本 PR，关闭 #39908