QServe：把 W4A8KV4 的理论收益做成实测吞吐的量化—系统协同设计

1. 摘要与一句话判断

3 分钟速读

1. 问题：现有 INT4 方案只加速低批量边缘推理，在大批量云端服务上失效——最先进的 W4A4 系统 Atom 在 A100 上跑 Llama-2-7B 比 TensorRT-LLM 的 W4A16/W8A8 慢 20-25%；已有 INT4 方法在 GPU 上去量化权重或部分和时承受 20-90% 运行时开销；A100 上一次 CUDA 核运算的代价相当于 50 次 INT4 张量核运算（第 1-2 页）P。
2. 为什么是 W4A8KV4：roofline 分析（A100 峰值 FP16/INT8/INT4 张量核 312/624/1248 TOPS、显存带宽 2TB/s）显示 GEMM 计算强度≈批量 m：m<78 时 W4A16（省带宽）更优、m>78 时 W8A8（算力高）更优，而 W4A8 只要全部计算落在 INT8 张量核上就两头通吃（第 3-4 页）P；解码注意力是计算强度 1 的 GEMV，KV4 等效加倍带宽（第 4 页）P。而 W4A4 的理论增益在 Ampere/Hopper 上兑现不了（第 4 页）P。
3. 算法（QoQ）：渐进分组量化——先 per-channel 对称 INT8（保护区间 [-119,119]，保证两级缩放反量化不溢出 INT8），再 per-group 非对称 INT4（组 128）；SmoothAttention 用 λ=max(|K|)^0.5 平滑 Key 缓存离群（约为其他激活的 ~10×），并按 RoPE 配对约束 λ_i=λ_{i+D/2} 融合进 W_Q/W_K（第 5-6 页）P。
4. 系统（QServe）：运行时精度映射为「所有 GEMM＝W4A8 输入、INT8 张量核计算、FP16 输出；所有注意力＝FP16 CUDA 核」（第 7-8 页图 11）P；计算感知权重重排按 32×32 tile 的「计算顺序」落盘、32 通道拼成单个 128bit 字，消除指针算术（第 8-9 页）P；解包 32 个 INT4 仅 3 次逻辑运算 + vadd4 单指令 4 路 INT8 加法实现寄存器级并行（第 9-10 页）P；KV4 注意力靠位技巧等四项优化避开 CUDA 核 roofline（第 10 页）P。
5. 结果与边界：同显存预算内最大可达吞吐（输入 1024/输出 512）对 TRT-LLM v0.9.0 最优精度配置：A100 8 模型 1.20×~2.38×、L40S 8 模型 1.39×~3.47×，几何平均 L40S 2.36×、A100 1.68×（第 13 页表 4、第 11 页图 15）P；34B 以下 7 模型中的 5 个在 L40S 上用 QServe 跑赢 A100 上的 TRT-LLM（第 13 页 §6.3）P；精度方面 WikiText2 困惑度 Llama-2-7B QoQ g128 5.67 vs FP16 5.47，五项常识任务平均对 FP16 仅损失 1.03%/0.89%/0.40%（7B/13B/70B，第 11-12 页）P；TTFT/TPOT、单请求延迟论文未报告I（口径核对）。

2. 问题背景

量化位宽与 GPU 硬件层次错配：P（第 1 页引言）现有整数量化分三类：W8A8、W4A16（前两类近似无损）、W4A4（精度受损）。INT4 理论上省显存、张量核峰值翻倍，但「降低位宽不一定加速推理」——关键在去量化落在低吞吐 CUDA 核上：A100 上一次 CUDA 核运算≈50 次 INT4 张量核运算（第 1-2、4 页）；现有 INT4 方法去量化权重或部分和的开销达运行时的 20-90%（第 1 页摘要）。

W4A4 系统的反例：P（第 1-2 页）在 A100 上服务 Llama-2-7B 时，Atom 比 TensorRT-LLM 的 W4A16/W8A8 慢 20-25%；图 2b（Llama-2-7B、A100、batch 64）印刷柱值：TRT-LLM-W8A8 2,104 tokens/s，而 Atom 986、QuaRot 817——尽管 INT4 张量核理论峰值是 INT8 的 2×（第 3 页图 2b 图注）。

roofline 解释（为什么 W4A8KV4 两头占优）：P（第 3 页图 3 与脚注、第 4 页 §3.1）A100 峰值 FP16/INT8/INT4 张量核 312/624/1248 TOPS、DRAM 带宽 2TB/s。GEMM 计算强度（MACs/element）≈m（序列数）：m<78 时 W4A16 理论吞吐更高（省权重带宽），m>78 时 W8A8 更优（INT8 算力）；W4A8 只要全部计算在 INT8 张量核上即可跨 batch 兼得。解码注意力是计算强度 1 MAC/元素的批量 GEMV，访存由 KV 缓存主导，KV4 等效加倍带宽、对 KV8 有 2× 理论峰值；batch=64 时注意力占总运行时 50% 以上（第 3-4 页）。

为什么不是 W4A4KV4：P（第 4 页 §3.2）张量核 GEMM 采用输出固定（output stationary）数据流，m×n×k 问题的主循环沿归约维 k 顺序迭代、超过 100 次迭代并主导运行时；W4A16 需在主循环内做 INT4→FP16 权重转换、Atom-W4A4 需 INT32→FP32 部分和转换；A100/H100 上 FP32 CUDA 核峰值仅为 INT4 张量核的 2%，Atom 去量化单个部分和≈50 次张量核 MAC，且为部分和维护 FP32+INT32 双份寄存器、限制可驻留 warp 数、削弱延迟隐藏。

目标工作负载与约束：P（第 3 页 §3.1）以真实对话负载为口径（输入 1024 token、输出 512 token），解码阶段运行时约为 prefill 的 6×，因此聚焦解码期 attention 与 GEMM；目标场景是大批量、云上 LLM 服务（第 1 页摘要：现有 INT4 只加速低批量边缘推理）。

基线口径：P（第 11-12 页 §6.1/§6.3）效率基线为 TensorRT-LLM v0.9.0（FP16 / W8A8 / W4A16 三种精度取最优）以及 Atom（W4A4）、QuaRot（W4A4，main 分支截至 2024-04-18）；精度基线为 FP16 及 W8A8/W4A16 量化方法（SmoothQuant、GPTQ-R、AWQ）。I 因此论文数字回答的是「QServe 比 2023-2024 年初的这些基线强多少」，不回答「与 2026 年的最新服务栈比谁优」。

3. 核心机制

3.1 QoQ 算法：为「去量化免费」而设计的量化格式

渐进分组量化与保护区间（§4.1）

P（第 5 页式 4/5、图 6）两级量化：先对权重做 per-channel 对称 INT8 量化（FP16 通道缩放），再对中间 8bit 张量做 per-group 非对称 INT4 量化（无符号 4bit 零点 + 无符号 8bit 组缩放，组大小 g=128 时记作 W4A8KV4g128）：

P（第 5 页「Protective Quantization Range」小节）直接级联会溢出：若某组 8bit 权重落在 [-113, 120]，INT4 非对称量化产生缩放 16、零点 7，值 120 反量化为 (15−7)×16=128，超出 INT8 上限 127；硬件饱和指令会拖慢至多 67%。推导：s_u8 ≤ (127−(−128))/(15−0)=17，要求 q̂_s8 ≤ 127 ⇒ q_s8 ≤ 119.5，故把第 1 级对称量化区间从 [-127,127] 收缩到保护区间 [-119, 119]，使反量化中间值严格落在 INT8 表示范围内（第 5-6 页）。

P（第 6 页「Compared to previous two-level quantization」）与 VSQuant、QLoRA 的 DoubleQuant、DGQ 的区别：它们先用目标位宽组量化、再对 FP16 组缩放做 per-channel 量化，GPU 上须先把缩放、再把权重反量化为浮点值，限制峰值吞吐；DGQ 服务系统把去量化 kernel 与 GEMM 分开，其 W4A8 GEMM 端到端延迟甚至慢于 cuBLAS 的 W8A8 GEMM。QoQ 的保护区间允许把去量化融合进 W4A8 GEMM kernel 并全程寄存器级并行——QServe 的 W4A8 per-group GEMM 对 W8A8 cuBLAS GEMM 达 1.5× 加速（GEMM kernel 级结论，非端到端吞吐）（第 6 页）P。

SmoothAttention（§4.2）

P（第 6 页图 7、式 7-9）经验观察：Value 缓存无明显离群，Key 缓存每头存在固定离群通道（幅值约为多数激活值的 ~10×），KV8 尚可容忍、KV4 量化级数不够。按通道因子 λ 平滑：Z=(QΛ)(KΛ⁻¹)ᵀ，λ_i=max(|K_i|)^α，实践 α=0.5 即可；因 Query 不量化只需处理 Key。为与 RoPE 交换律兼容（RoPE 把通道 i 与 i+D/2 配对），加硬约束 λ_i=λ_{i+D/2}（式 9）；平滑因子 Λ 融合进前层权重（W_Q=ΛW_Q、W_K=Λ⁻¹W_K），运行时零额外 kernel 调用（第 6 页；「无系统开销」结论见第 13-14 页图 16）P。

通用量化优化（§4.3）

• 块输入模块旋转（§4.3.1）：P QKV 投影与 FFN 第 1 层等「输入模块」前施加 scaled Hadamard 旋转；旋转是酉变换，可被前一块输出模块权重吸收（第 7 页）。
• 块输出模块平滑（§4.3.2）：P out_proj 与 FFN 第 2 层等「输出模块」用对角 Λ 平滑中间激活并吸收进前层权重；若照搬输入模块的迁移强度，Llama-2-7B 的 WikiText-2 困惑度会至多恶化 0.05，故迁移强度 α 取近 0、λ 基本由权重决定（第 7 页）。
• 激活感知通道重排（§4.3.3）：P 用 max(|X|) 定通道显著性，ArgSort 使相近显著度通道落入同组，替代 Atom 的混合精度方案（第 7 页图 10）。
• 权重裁剪（§4.3.4）：P 多数线性层最小化层输出误差，q_proj/k_proj 例外、最小化块输出 MSE（第 7 页式 10）。

3.2 QServe 服务系统

运行时精度映射（§5.1，图 11）

P（第 7-8 页）所有 GEMM 层以 W4A8 为输入、在 INT8 张量核上计算、输出 FP16；所有注意力层在 CUDA 核上以 FP16 计算；每个 LLM 块 FP16 进、FP16 出。激活量化融合进前趋 LayerNorm（QKV 投影、FFN 第 1 层）或前趋激活 kernel（FFN 第 2 层）；注意力 out_proj 前插入独立量化节点。

KV 缓存管理（§5.1）

P（第 8 页）沿用 vLLM / TensorRT-LLM 的分页 KV 缓存防碎片，并支持同类 in-flight batching；但两框架是 per-tensor 静态（离线算好缩放）的 KV8 量化，QServe 因 4bit 精度需要 per-head 动态 KV 量化——FP16 缩放与零点按头紧跟量化 KV 特征存于每个缓存页内、可在线更新。

计算感知权重重排（§5.2.1，图 12）

P（第 8-9 页）朴素 INT4 寻址每 4 通道要一次地址计算，而 CUDA 核吞吐仅为 A100 INT8 张量核的 32×，指针算术不可忽略；跨步访问还无法打包 128bit 读。ldmatrix 指令只在存储/计算位宽一致（W8A8）时可用，W4A8 按「字节一致、元素不一致」分配而失效。解法：按 32×32 tile 把计算所需顺序落盘（线程 0 用输入通道 0-3 与 16-19 对应的权重，32 通道拼成单个 128bit 字），权重静态、零运行时开销，同时消除指针算术并保证 128bit/thread 访存；零点与缩放同样重排。

快速去量化（§5.2.2-5.2.3，图 13-14）

P（第 9-10 页）把 ZINT4→SINT8 转换分解为 UINT4→UINT8 解包与 UINT8→SINT8 零点减两步：

• 解包：P 32 个 UINT4 权重重排为 w0,w16,w1,w17,…，用寄存器级并行仅 3 次逻辑运算解包为 UINT8（第 9 页图 13）。
• per-channel 场景：P 零点减法移出主循环并入 GEMM 尾注——X(Z_S_W)=t_X×(z_W⊙s_W)，其中 t_X=X·1_k 预计算；每个 W4A8 kernel 前必是访存受限 kernel，预计算可融合进去、时延开销可忽略（第 9 页式 13）。
• per-group 场景：P 零点按组、无法并入尾注，且二级缩放引入逐权重 INT8 乘法；采用「先乘后减」次序配合 vadd4（一条 INT32 ALU 完成 4 路 INT8 加法）实现 4 路寄存器级并行——GPU 没有 4 路 INT8 乘法指令，靠 8bit 缩放高位补 24 个零模拟，仅当乘积不溢出 INT8 时成立，渐进量化的保护区间正保证这一点（第 9-10 页、第 10 页开头）。

通用 GEMM 优化（§5.2.4）与 KV4 注意力（§5.3）

P（第 10 页）访存侧：多级软件流水 + 异步内存拷贝、L1 共享内存 swizzle 消除 bank 冲突、跨线程块置换计算分区提升 L2 复用（相邻块复用同一权重）；计算侧：m 较小时把归约维 k 切片、在 L1 内跨 warp 归约部分和。KV4 注意力部分与实验数字见 #experiments（表 1）。

4. GPU/系统数据路径

端到端文字序列（与图中编号一致）

1. ① 服务层：请求进入引擎调度器做 in-flight batching（与 vLLM / TensorRT-LLM 同类）；分页 KV 管理器按页分配显存防碎片，每页内联存放该头 FP16 缩放/零点（KV4 专用、在线更新）P（第 8 页 §5.1）。
2. ② 块内精度映射：块级 FP16 进、FP16 出；LayerNorm（QKV 投影、FFN 第 1 层前）或激活 kernel（FFN 第 2 层前）融合激活量化输出 INT8；QKV 投影、FFN 两层均为 W4A8 GEMM（INT8 张量核、FP16 输出）；注意力在 CUDA 核上 FP16 计算并读 KV4；out_proj 前插入独立量化节点P（第 7-8 页图 11）。
3. ③ W4A8 GEMM 主循环：加载按 32×32 tile「计算顺序」离线重排的 INT4 权重、128bit 连续读 → 3 次逻辑运算解包为 UINT8 → 乘后减次序 + vadd4（一条 INT32 ALU＝4 路 INT8 加法；保护区间 [-119,119] 保证乘积不溢出）→ INT8×INT8 张量核矩阵乘 → 尾注处理（per-channel 零点并入尾注、t_X 预计算融合进前趋访存受限 kernel；per-group 二级缩放在主循环内寄存器级相乘）→ 输出 FP16P（第 5、8-10 页图 5d、12-14）。
4. ④ 解码注意力：分页读 INT4 K/V，位技巧把反量化算术强度降到每元素 2 次运算，QKᵀ 与 SV 积用 FP16，kernel 起始异步预取缩放/零点；FP16 softmax 与输出。风险：朴素 KV4 每元素 5 次 ALU、A100 FP32 CUDA 核 roofline 转折点仅 9.8 Ops/Byte，融合核会变成计算受限（A100 上 1.1-1.2× 反而变慢）；QServe 四项优化后对 TRT-LLM KV8 核达 1.5×（kernel 级）P（第 10 页 §5.3、表 1；1.7× 为端到端注意力改进，第 14 页 §6.4）。
5. ⑤ KV 写路径：新 token 的 k/v 逐头做非对称 INT4 动态量化（缩放/零点在线更新），INT4 特征与该头 FP16 缩放/零点内联写入缓存页；SmoothAttention 已离线融合进 W_Q/W_K，运行时零额外调用。每步解码循环 ②→⑤P（第 8 页 §5.1、第 6 页 §4.2）。

5. 架构权衡

• 精度组合选择 ↔ 精度代价

  P W4A8KV4 不是免费的：Llama-2-7B WikiText2 困惑度 FP16 5.47 → QoQ(W4A8KV4) 5.75（per-channel）/ 5.67（g128）；对 W8A8 SmoothQuant 与 W4A16 AWQ 至多增加 0.16；五项常识任务平均对 FP16 损失 1.03%/0.89%/0.40%（7B/13B/70B，第 11-12 页表 2/表 3）。KV4 这一步在消融中再恶化 0.14（5.66→5.80，第 13-14 页图 16）。I 即 QServe 用「可量化的精度预算」换 INT8 张量核路径，接受与否取决于业务对困惑度/下游任务的敏感度。

• 理论 roofline ↔ 实测可实现性

  P KV4 理论上对 KV8 有 2× 峰值，但朴素替换在 A100 上 1.1-1.2× 反而变慢（CUDA 核 roofline 转折 9.8 Ops/Byte 提前到来，第 10 页表 1 及 §5.3）；W4A4 理论上 m>78 占优，但 Atom/QuaRot 实测显著落后（第 3-4 页）。I 论文的核心方法论即「位宽选择必须与 kernel 实现协同」——引用任何「INT4 更快」的推导前先看实现是否避开了 CUDA 核瓶颈。

• 算法—系统协同 ↔ 设计耦合

  P 渐进量化的保护区间是为「乘后减 + vadd4 的寄存器级并行」能安全工作而反推出来的量化格式约束（第 5-6、9-10 页）；SmoothAttention 以 RoPE 配对硬约束换取融合进权重、运行时零开销（第 6 页）；权重重排按 kernel 瓦片顺序落盘（第 8-9 页）。I 三者都把算法格式与具体 GPU 指令（vadd4、128bit 访存、ldmatrix 的局限）绑定——换硬件架构（不同 ALU/访存粒度）时这套协同需要重新验证，不是纯算法层可移植的成果。

• 吞吐收益 ↔ 延迟口径缺失

  P 论文全部效率结论是「同显存预算内最大可达吞吐」（batch 由显存约束扫描决定），未报告 TTFT/TPOT/单请求延迟（第 12 页 §6.3；全文无延迟 SLO 指标）。同 batch 拆解显示 1.88×（L40S、Llama-2-7B）＝1.45×（同 batch）×1.3×（更大 batch），即约三分之一收益来自「KV4 省出的显存装了更大 batch」（第 14 页图 17 及正文）。I 以单请求延迟为首要 SLO 的场景不能直接引用这些倍数。

• 显存容量 ↔ 基础设施降本

  P W4 权重量化在消融中省 3.5GB、KV4 使 GPU 内存使用减半（第 13-14 页图 16 文字结论；显存面板无数值标注）；L40S-48G 因此能容纳与 A100-80G 相同的 batch，34B 以下 7 模型中 5 个「L40S+QServe」吞吐超过「A100+TRT-LLM」，表 4 表注据此称「有效把 LLM 服务美元成本降低 3×」（第 12-13 页）。I 该 3× 是论文级断言，其隐含的 L40S/A100 价格比未在论文中披露，落到具体云价需另行核算（见 #cost-slo）。

• 实测领先 ↔ 基线陈旧

  P 效率基线为 TRT-LLM v0.9.0（2023 年版本）与 QuaRot/Atom 的 2024-04-18 main 分支（第 11 页 §6.1）；论文附录自述「绝对吞吐因机器条件略有差异，相对加速比应保持一致」（第 18 页表 6 注）。I 对照当下软件栈（TRT-LLM 已多年迭代，且官方仓库 README 称 QServe 已于 2024 年 12 月集成进 TensorRT-LLM，R），绝对数字不可代表当前版本表现；决策前应在目标栈上复测，论文数字用于理解机制与量级。

6. 云上部署映射

以下映射为厂商中立示例；具体产品命名/规格仅作说明并标 E，以厂商当时目录为准，未逐一在线核验。论文事实单独标注 P。

7. 成本 / 性能 / SLO

7.1 指标口径（论文使用的）

P（第 12 页 §6.3）效率主指标＝同显存约束内最大可达吞吐（tokens/s；输入 1024 / 输出 512；A100 80GB / L40S 48GB 预算；A100 用 per-channel、L40S 用 per-group 权重量化）。精度指标＝WikiText2 困惑度（序列 2048，越低越好）、五项常识任务零样本平均（PIQA/ARC-e/ARC-c/HellaSwag/WinoGrande，越高越好）、LongBench 长上下文。TTFT、TPOT、单请求延迟、能耗：论文未明确披露。

7.2 主效率结果上下文（表 4 全量转录，第 13 页）

P（第 11 页图 15、第 13 页表 4、第 12-13 页 §6.3）补充口径：图 15 几何平均加速比对 TRT-LLM v0.9.0 为 L40S 2.36×、A100 1.68×（同为 80GB/48GB 显存预算）；表 4 表注「QServe 在 L40S 上对比 TRT-LLM 在 A100 有竞争力吞吐，有效把 LLM 服务美元成本降低 3×」；§6.3 正文「34B 以下 7 个模型中的 5 个在 L40S 用 QServe 服务甚至高于 A100 用 TRT-LLM」「L40S 上跑 Llama-2-7B 比 Atom 高 10% 吞吐」；对 Atom/QuaRot（A100）为 2.5-2.9×（第 2 页引言）。I 摘要「1.2×/1.4×/2.4×/3.5×」与表 4「1.20×/1.39×/2.38×/3.47×」为舍入关系；§6.3 文字「1.47× 到 3.47× 覆盖七个模型」与表 4 八列（最小 1.39×）存在口径出入——页面以表 4 精确值为准并保留各处原始表述。

7.3 容量与成本公式（参数化，编辑推断）

I 公式中可由论文支撑的量：KV4 使 KV 显存约减半、W4 省 3.5GB（Llama-2-7B，图 16 文字结论）、同预算内吞吐倍数（表 4）；其余（单价、激活开销、回退余量、请求长度分布）需按部署实测填写，本页不给出伪精确数字。P 论文对「美元成本降 3×」的表述以 L40S 吞吐 ≥ A100 吞吐为前提（表 4 表注），未给出任何价格数据。E 云单价随时段/区域/折扣变化，核算时以查询当时厂商目录为准并记录日期。

7.4 敏感项与数据缺口

• P 敏感项（各自条件下）：模型规模（越小加速越保守——Llama-3-8B 在 A100 仅 1.20×，72B 级 2.38×-3.47×，表 4）；GPU 代际（L40S CUDA 核更强、收益整体高于 A100，第 12 页 §6.3）；量化粒度（A100 per-channel vs L40S per-group，第 12 页）；batch（同 batch 加速 1.45× 与更大 batch 因子 1.3× 叠加，第 14 页图 17）。
• P 数据缺口：TTFT/TPOT/单请求延迟、能耗、不同上下文长度下的吞吐曲线、校准数据集（默认 Pile validation）之外的超参敏感性、KV 页内统计量的空间占比——论文未明确披露。
• I 云价格、量化流水线（校准+离线重排）的一次性工程成本、每轮模型升级的再量化成本：论文未涉及，部署前自测。

8. 安全与可运维性

8.1 安全

• KV 缓存即用户数据：I KV4 量化不改变 KV 缓存的语义——它仍逐 token 保留对话状态（只是精度降为 4bit 并内联 per-head 统计量），量化不是加密。多租户实例复用前必须重置/清零 KV 显存（含页内缩放/零点），显存 dump 即数据泄露。
• 数据驻留与边界：P KV 缓存驻留 GPU 显存（分页管理，第 8 页），论文无外置/卸载设计。I 云上驻留边界＝实例边界；跨可用区不复制 KV。
• 供应链：R 代码来自摘要印明的官方仓库 github.com/mit-han-lab/omniserve（第 1 页；2026-09-15 在线核验：Apache-2.0、main 分支、未归档）。P 论文附录 A 提供预构建 Docker 镜像（shang12138/qserve-mlsys25-ae）、代码许可 Apache License 2.0、数据许可 MIT、制品存档 DOI 10.5281/zenodo.14991385（第 18 页）。I 引入时 pin commit/镜像 digest、扫描 CUDA/PTX kernel 与依赖、保留可回退版本；第三方镜像内容与论文制品的一致性需自行校验。
• 输出风险继承：P 量化改变输出分布：WikiText2 困惑度与常识/长上下文基准的平均退化见 #experiments（第 11-13 页）。I 合规审查应以「量化后模型的实际输出」为对象；长尾任务退化可能大于均值口径。

8.2 可运维性（含恢复与回滚）

• 回退路径：I 回退＝切换到基线服务栈（论文口径即 TRT-LLM FP16/W8A8/W4A16）。关键运维约束：回退后 KV 显存占用上升（KV4→KV8 翻倍）、同容量下 batch 变小，必须预留显存余量或同步降载，否则回退动作本身会触发 OOM（表 4 中多款大模型基线即 OOM）。
• 配置变更管理：P 量化粒度与 GPU 绑定（A100 per-channel / L40S per-group）、组大小 g=128 与否是两个独立格式（W4A8KV4 vs W4A8KV4g128，第 10-11 页）。I 把「量化格式 + GPU 型号 + 模型版本」作为绑定配置资产管理，变更走灰度 + 精度回归，禁止热改。
• 离线量化流水线运维：P 每个模型上线前需跑校准（默认 Pile validation）→ 旋转/平滑/重排/裁剪 → 两级量化 → 权重重排落盘（第 5-7 页、第 12 页 §6.2）。I 这是一次性但必须纳入发布流程的环节：新模型、新版本都要重量化 + 重评精度；建议归档校准集快照与量化超参保证可重建。
• 故障恢复：I 推理服务无训练态，故障恢复＝实例重启 + 模型/权重制品重载；KV 缓存不持久化、丢失即重算（prefill 重放）。多实例无状态水平扩展，单实例故障由负载均衡摘除（云编排能力，E）。
• 监控告警：I 建议：①吞吐相对验收基线的漂移（kernel 性能回归或调度异常）；②显存水位与 OOM 边界批量（KV 页内统计量异常增长或碎片化会先表现为容量下降）；③困惑度/下游任务抽样分数（量化格式配错时先伤精度）；④in-flight batching 队列深度与 token 延迟分位数（论文未提供延迟口径，需自建基线）。
• 升级与回滚：I CUDA/PTX kernel 与框架版本配对 pin；升级走金丝雀 + 双基线回归（开/关量化各测吞吐与精度）；回滚＝上一配对制品或切回基线栈，两者独立可执行。复现环境参考附录 A：x86_64 主机 + 至少一块 A100（推荐）或 L40S、磁盘约 512G、拉镜像约 1 小时、QServe 基准约 1 小时、TRT-LLM 基准 2-4 GPU 小时P（第 18 页）。
• 验收基线：I 上线前固定「基线栈 vs QServe」两组读数：同预算最大吞吐、OOM 边界批量、困惑度与硬任务抽样，形成可复核差值；此后任何变更以差值回归验收。附录 A 提醒：绝对吞吐因机器条件略有差异，相对加速比应保持一致P（第 18 页表 6 注）。

9. 适用 / 不适用场景

适用（触发条件 + 理由）

1. 大批量、成本敏感的云端推理服务，且显存预算限制并发。触发：同显存预算下吞吐不达 SLO，或批量被显存卡住。理由：同预算内最大可达吞吐对 TRT-LLM v0.9.0 在 A100 提升 1.20×~2.38×、L40S 提升 1.39×~3.47×（8 模型，输入 1024/输出 512）P（第 13 页表 4）；KV4 让 L40S-48G 容纳与 A100-80G 相同 batchP（第 13 页 §6.3）。
2. 采购/扩容决策点上的「低价位卡 + 量化」路线评估。触发：新购 GPU 或云实例选型。理由：论文显示 34B 以下 7 模型中 5 个「L40S+QServe」吞吐超过「A100+TRT-LLM」，表 4 表注称有效降低美元成本 3×P（第 12-13 页）；I 落地前须以当期云价与自身流量复算。
3. KV 显存压力大、上下文较长的批量服务。触发：KV 占显存比例高、想在不换卡的情况下扩 batch 或上下文。理由：消融显示 KV4 使 GPU 内存使用减半、该步吞吐 +1.47×（L40S、Llama-2-7B、64 请求×1024/512）P（第 13-14 页图 16）；长上下文精度退化小（LongBench 平均 38.52→38.38，Llama-3.1-8B-Instruct）P（第 13 页表 5）。
4. 需要「INT8 张量核友好」量化格式的自研 kernel 团队。触发：自建推理栈、想复用 W4A8 主循环设计。理由：保护区间 + 乘后减 + vadd4 的 RLP 设计、计算感知权重重排均有完整推导与消融P（第 5-10 页）；官方仓库 Apache-2.0 允许商用集成R（2026-09-15 核验）。

不适用（触发条件 + 理由）

1. 低批量、边缘或端侧部署。触发：batch 小（GEMM 计算强度 m<78）、显存富余。理由：论文明确「现有 INT4 只加速低批量边缘推理」的反面诉求是大批量云端服务，且其 roofline 显示小 m 时 W4A16 的省带宽路径理论上已更优P（第 1 页摘要、第 3 页 §3.1）；收益对象是批量吞吐而非单请求。
2. 以单请求延迟（TTFT/TPOT）为首要 SLO 的交互式场景。触发：目标是个位数毫秒级每 token 延迟或首字延迟。理由：论文只报同预算最大可达吞吐，未报告 TTFT/TPOT/单请求延迟P（第 12 页 §6.3；全文无延迟指标）。I 「吞吐倍数」不能自动兑换成「延迟倍数」。
3. 追求 W4A4 理论峰值、且有能力改造 kernel 的极端显存压缩场景。触发：需要 W4A4 才能装下的模型/批量。理由：QServe 的立场是 W4A4 理论增益在现有 GPU 架构上兑现不了、其实现走 INT8 张量核路径，不提供 W4A4 主循环P（第 3-5 页 §3.2、图 5）。
4. 把官方仓库当活跃开源产品依赖的团队。触发：需要长期社区支持、多框架适配、新模型开箱即用。理由：仓库截至 2026-09-15 共 25 commits、无 release、README 最新动态为 2025 年 2 月（「OmniServe is now publicly available!」），近一年半近乎无更新R（2026-09-15 核验）；README 称 QServe 已于 2024 年 12 月集成进 TensorRT-LLM，I 生产路径建议优先评估框架内置实现而非自行维护论文仓库。
5. 对量化精度零容忍的业务。触发：无法接受任何困惑度/下游任务退化。理由：QoQ 相对 W8A8/W4A16 至多增加 0.16 困惑度、对 FP16 有 0.40%~1.03% 常识任务损失、KV4 一步 +0.14 困惑度P（第 11-12、13-14 页）——非零退化是设计前提。

10. 实验与指标

10.1 实验设置与口径 P（第 10-12 页）

10.2 精度结果（表 2 关键列 / 表 3 全量 / 表 5）

P（对比表述在第 12 页 §6.2 正文；表 2 数值在第 11 页）口径：对 Llama-2-7B，QoQ 相对 W8A8 SmoothQuant 与 W4A16 AWQ 至多增加 0.16 困惑度，相对 W4A4 QuaRot 至多改善 0.49。表 2 共 10 个模型列（本页转录 Llama-2-7B 全配置对比与 QoQ g128 vs FP16 的多模型对比；完整逐格值见 PDF 第 11 页）。

P（结论句在第 12 页 §6.2 正文；表 5 数值在第 13 页）结论：QoQ 相对 BF16 基线在长上下文上保持最小退化（平均 38.52→38.38）。注：表 5 表头为斜排文字，文本提取易错位；本表列序与数值已按 220dpi 放大原页逐列复核（DuReader Retrieval 至 MultiFieldQA-En 从左至右），纠正了早期提取中任务名与数值的一列错位。

10.3 效率与消融

KV4 注意力 kernel（表 1，A100，核级延迟）

P（第 10 页）表 1 图注：朴素 KV4 注意力在 L40S 上比 TRT-LLM-KV8 快 1.7×，但在 A100 上因 CUDA 核 roofline 转折点提前而慢 1.1-1.2×。四项优化：FP32 全换 FP16（计算 roof 翻倍）、位技巧把反量化算术强度降到每元素 2 次运算、简化控制逻辑、kernel 起始预取缩放/零点——对 TRT-LLM KV8 核达 1.5×（A100，attention kernel 级）（第 10 页 §5.3）。P（第 14 页 §6.4）改进拆解（64×1024 输入）：基础 KV4 0.48ms → 位技巧 0.44 → 简化控制流 −0.05 → QKᵀ/SV 转 FP16 各 −0.03 → 预取后 0.28ms，端到端注意力改进 1.7×。I 注意区分三个口径：表 1 是注意力核延迟；1.7× 是注意力环节的端到端改进；表 4 的 1.20×~3.47× 才是整服务吞吐。

量化技术消融（图 16，Llama-2-7B、L40S、64 请求×1024/512）

I 图 16 的 GPU 显存面板（权重 + KV 缓存堆叠柱，坐标 0-32GB）无数值标注，论文文字仅承诺「W4 省 3.5GB」「KV4 使 GPU 内存使用减半」两处结论；任何具体 GB 读数（如全配置约 16GB）都是编辑读图近似，本页不引用。

系统消融与去量化开销（图 17-18，第 14 页）

P（第 14 页）同 batch 对比（L40S，输入 1024/输出 512）：Llama-2-7B 对 TRT-LLM 的 1.88× 可分解为 1.45×（同 batch）×1.3×（更大 batch）；Llama-2-13B 两个因子同等重要。图 18：QServe per-group W4A8 GEMM 的去量化开销与 TRT-LLM-W4A16 相当、远小于 Atom-W4A4（后者至多 90%）——但 QServe 计算在 INT8 张量核上、享受 2× 更高吞吐。I 图 18 为柱状图、无数值标注，上述为图注/正文文字结论，不代表逐 batch 精确占比。

10.4 建议复现步骤（编辑推断）

1. I 取官方仓库代码（已核验归属/许可，pin commit/digest）；或直接使用附录 A 的预构建 Docker 镜像（x86_64 + A100/L40S、约 512G 磁盘）P（第 18 页）。
2. I 精度复测：目标模型跑 FP16/BF16 基线与 QoQ W4A8KV4g128（校准集默认 Pile validation），对齐 WikiText2（序列 2048）、五项常识任务（lm-eval）与 LongBench 口径；可用开源库 deepcompressor 复现精度结果（论文仅提名称，未给 URL）P（第 11-13、18 页）。
3. I 效率复测：A100-80G-SXM4 与 L40S-48G、输入 1024/输出 512，从 batch 2 起扫描至显存预算上限，记录最大可达吞吐；对照表 4 形状（小模型差距小、大模型/低配基线 OOM 处差距大）。附录 A 提醒绝对吞吐随机器条件浮动、应以相对加速比为准P（第 18 页）。
4. I 目标业务实测：以真实上下文长度分布与并发区间测「基线栈 vs QServe」差值（吞吐、显存、硬任务抽样精度），并补测论文缺失的 TTFT/TPOT；再决定启用范围与量化粒度（per-channel/per-group）。
5. I 若生产栈为 TensorRT-LLM：先评估框架内置实现（官方仓库 README 称 2024 年 12 月已集成R），再决定是否需要论文仓库级自建。

11. 论文 / 代码 / 延伸链接

12. 给架构师的决策清单

I 以下为落地前的勾选项；标注（P）的条目对应论文证据，（R）对应仓库核验项，其余为工程判断。

• 需求与规模
  • 确认瓶颈与收益对象匹配：大批量、同预算吞吐不足或 batch 被显存卡住，而非单请求延迟问题（P：收益口径＝同显存预算内最大可达吞吐，第 12 页 §6.3；TTFT/TPOT 论文未披露）
  • 明确批量区间：小批量（GEMM m<78）时论文的 roofline 论证不倾向 W4A8（P：第 3-4 页）
  • 盘点模型清单与量化格式映射：W4A8KV4g128 vs W4A8KV4（per-group vs per-channel）按 GPU 选型（P：A100 per-channel / L40S per-group，第 12 页）
• 兼容性
  • 优先评估生产推理框架的内置实现/等效配方（官方 README 称 2024-12 已集成进 TensorRT-LLM）（R：2026-09-15 核验；I：选型判断）
  • 若自建：确认目标 GPU 的 vadd4/128bit 访存等指令前提与 kernel 移植成本（P：RLP 设计绑定 CUDA/PTX 与 A100/L40S，第 9-10 页；I：移植评估）
  • 确认与现有系统组件（分页 KV、in-flight batching、投机解码等）的叠加关系，论文未测这些组合（P：第 8 页仅自述同类；I：叠加需自测）
• PoC（1-2 周量级）
  • 复测「基线栈 vs QServe」同预算最大可达吞吐（输入 1024/输出 512 起步，再换真实分布）（P：第 12-13 页；I：方法）
  • 精度回归：WikiText2 困惑度 + 2-3 个业务下游任务，对照论文容忍度（QoQ vs W8A8/W4A16 至多 +0.16 困惑度；KV4 一步 +0.14）（P：第 11-12、13-14 页）
  • 补测论文缺失口径：TTFT/TPOT 分位数、不同上下文长度吞吐曲线（P：论文未披露；I：自建基线）
  • 记录 OOM 边界批量，验证「L40S 容纳 A100 同 batch」的容量结论在自身负载下成立（P：第 13 页 §6.3；I：实测要求）
• 容量
  • 按 §7.3 公式重算单实例可容 batch：显存预算 − W4 权重（含两级统计量）− KV4（含 per-head 统计量）− 激活/框架开销 − 回退余量（P：页结构，第 8 页；「W4 省 3.5GB、KV4 减半」，第 13-14 页；I：公式与余量）
  • 为回退路径预留显存：KV4→KV8 显存翻倍、同容量 batch 变小，回退动作本身不得触发 OOM（I：回退膨胀推算；P：表 4 基线 OOM 分布，第 13 页）
  • 72B 级模型单独规划：48GB 预算下吞吐绝对值低（表 4 L40S 上 70B 仅 286 tokens/s），按业务量决定是否值得（P：第 13 页表 4）
• SLO 与恢复
  • 定义吞吐 SLO 并以「开/关量化」双基线固化验收差值（I）
  • 定义精度 SLO：困惑度/下游任务阈值，跌破即回退（P：精度代价非零是设计前提，第 11-14 页；I：阈值自定）
  • 故障演练：实例重启重载、切回基线栈降级各演练一次（I）
• 成本
  • 以「当期云价 + 查询日期」复算「L40S+QServe vs A100+基线」的单 token 成本，不把论文「降 3×」当价格结论（P：表 4 表注为论文断言、无价格数据，第 13 页；I：核算方法）
  • 计入一次性成本：离线量化流水线、每轮模型升级的再量化与再评测（P：附录 A 基准时长参考，第 18 页；I：成本项）
  • 用同 batch 拆解区分「kernel 加速」与「容量换 batch」两个收益来源，避免重复计价（P：1.88×＝1.45××1.3×，第 14 页图 17）
• 安全
  • KV 显存（含页内 per-head 缩放/零点）按用户数据处理：租户/会话切换清零，显存 dump 纳入威胁模型（P：页结构，第 8 页；I：残留风险）
  • 供应链 pin：官方仓库 commit/镜像 digest 锁定 + CUDA/PTX kernel 扫描 + 可回退层（R：Apache-2.0 已核验；P：附录 A 制品清单，第 18 页；I：流程要求）
  • 输出合规以「量化后模型实际输出」为审查对象（I：量化改变输出分布）
• 运维与回滚
  • 运行时「切回基线量化配置」开关就绪，并与降载联动、有告警（I）
  • 监控：吞吐漂移、显存水位/OOM 边界批量、困惑度与硬任务抽测、in-flight 队列深度与延迟分位数（I）
  • 「量化格式 + GPU 型号 + 模型版本」绑定变更管理：灰度 + 精度回归，禁止热改（P：粒度按 GPU 选型，第 12 页；I：治理）
  • 升级走金丝雀 + 双基线回归；复现环境对齐附录 A（Docker、x86_64、A100/L40S、约 512G 磁盘）（P：第 18 页；I：流程）
• 退出策略
  • 退出成本评估：切回基线栈即回退，无数据格式迁移；但容量回到 KV8/FP16 水位，需重算容量与预算（I；P：基线即回退态）
  • 归档量化制品、校准集快照、超参与双基线读数，保证可重建（P：附录 A 制品化思路，第 18 页；I：归档要求）
  • 若转向其他量化/缓存方案（KV8 服务栈、W4A16 weight-only 等），把「QServe 的 INT8 张量核路径设计」作为对比基线而非二选一（I：对比框架；P：图 5 各格式主循环对比，第 5 页）

13. 证据台账

下表为核心结论的证据映射；逐条引文与核验记录见构建文件 sources/qserve.evidence.json（31+ 条，EV-01~EV-32）。核验日期为 2026-09-15。