企业官网建设流程全解析

2080 Ti 跑 Qwen3.8-27B：给 vLLM 写了个插件，decode 快 1.6 倍，首字快 30%

两张 2080 Ti 22G（魔改大显存版）跑 Qwen3.8-27B-AWQ + MTP 投机解码，vLLM 默认配置下每秒只有 69 个 token 左右，而且不管上下文是 1K 还是 64K，每步都稳定在 47~51 ms —— 这个「不随上下文变」的时间，一看就不是 GPU 在忙。。

顺着这条线索挖下去，发现 vLLM 在老卡上有三处慢，于是写了个插件vllm-boost把它们换掉。不改 vLLM 源码、不改启动命令，pip 装上就生效：

• decode 吞吐1.558×（69 → 107 tokens/s）
• 首字时间−24~30%
• GSM8K 前 40 题的最终答案和原版完全相同

项目地址：https://github.com/Miaaanz/vllm-boost

先看结果

2 × RTX 2080 Ti 22G，Qwen3.8-27B-AWQ-INT4，MTP 每步猜 3 个 token，单请求。基线是 vLLM 0.27.1 默认配置，两组在同一台机器上用同一套脚本测。

长上下文也能用：约 25 万 token 的大海捞针答对，decode 每步约 44 ms。代价是 KV cache 容量少了约 3.8%。

vLLM 在老卡上慢在哪

2080 Ti、T4 这类 Turing 卡（算力 7.5）便宜、显存够，但 vLLM 和它依赖的库基本都围着新卡优化。跑 Qwen3.8 + 投机解码时，有三个问题叠在一起，单看 profile 很难发现：

1. 完整 CUDA Graph 被悄悄降级。老卡上 vLLM 默认用 FlashInfer 注意力后端，但它不支持把投机解码的验证步（每个请求 1 + 3 个 query）放进完整 CUDA Graph，vLLM 就自动降成分段图：每步几百个 kernel 由 CPU 一个个发射。这就是那 47~51 ms 的来源 —— 瓶颈在 CPU，和上下文长短无关。

2. 换成 Triton 后端也不行。Triton 后端能整图，但 Triton 从 3.3 起在算力 < 8.0 的卡上不生成张量核指令（triton#5066），注意力只有 0.7~1 TFLOPS。8K 上下文单层注意力要 598 µs，上下文越长越慢。

3. GDN 线性注意力也是 Triton 写的。Qwen3.8 的 64 层里有 48 层是 GDN 线性注意力。vLLM 只在 H100 / B200 这类卡上给它用 FlashInfer，其他卡都走 Triton 写的 flash-linear-attention，在 2080 Ti 上同样没有张量核，一层 4000 token 要 19 ms，48 层加起来约占首字时间的 1/3。

三个优化

1. 注意力：用 Triton 后端的整图能力，计算换成 FlashInfer

思路是两头取长处：留在能整图的 Triton 后端，只把注意力计算换成有张量核的 FlashInfer kernel。单层注意力（8K 上下文）598 µs →72 µs，prefill 从 0.7~1 TFLOPS 提到 24~35 TFLOPS。

麻烦在于让 FlashInfer 的 decode 进 CUDA Graph：

• plan 不能在图里做。FlashInfer 的 kernel 要先在 CPU 上 plan，而图重放时不跑 Python。所以每一步都在图外 plan，页表和调度信息写进固定地址的缓冲区，图里只捕获 run。
• CPU 上没有准确的序列长度。开投机解码时 vLLM 默认异步调度，CPU 只知道长度上限（假设草稿全被接受）。做法是按上限 plan，再在 GPU 上把最后一页的有效长度改成准确值。
• KV 块太大，切不开。混合模型的 KV 块被 vLLM 调成 800 token，FlashInfer 按页切分 KV，短上下文时并行度上不去。块内位置是等距的，所以不拷贝数据，直接把每块看成 50 个 16 token 的虚拟页。

2. GDN prefill：手写 CUDA 张量核，单层快 14 倍

算法和 flash-linear-attention 一样（块长 64），但用 CUDA 重写，用的是 Turing 就有的 mma.sync.m16n8k8 张量核指令，拆成三个 kernel：块内准备、块间状态递推、输出。块间递推原来占 72% 的时间，现在状态一直留在寄存器里（就是 mma 的累加器），下一块的输入在算这一块时就预取好。

一层 4000 token：19.19 ms →1.35 ms（14.2×），误差和原实现在同一量级。首字时间的提升主要来自这里。

3. 投机解码草稿：只在高频词里猜

这一项和硬件无关。MTP 草稿每猜一个 token 都要算一遍 lm_head 再取 argmax，Qwen3.8 词表 24.8 万，27B 的 lm_head 有 2.5 GB，每步猜 3 个就要读 3 遍，约占每步时间的 1/5。

但草稿只需要 argmax，而且猜错的 token 会在验证时被主模型拒绝、换成正确的 —— 所以只在最常用的65,536 个 token上算，最终输出和原来完全一样，每次读的权重从 1.27 GB 降到 0.34 GB（每张卡）。

两个踩过的坑：

• 双卡不均衡。vLLM 把 lm_head 按 token id 连续切给两张卡，而高频 token 的 id 大多很小（BPE 先合并的常用片段），直接取子集会有 63,353 个落在第 0 张卡、只有 2,183 个在第 1 张卡，第 0 张卡几乎没少读。所以加载时把子集按位置轮流分给两张卡，各算 32,768 行。
• 语言比例很重要。第一版词表用英文为主的文本统计，中文回答只覆盖 80.7%，平均每 5 个 token 就有一个草稿不可能猜中，中文对话的每步产出比不裁剪还低 8%，省下的读权重时间几乎被抵消（吞吐只高 3%）。按中英文重新统计后覆盖率到 98.7%，每步产出和不裁剪持平，中文对话吞吐比不裁剪高14%。

新卡上也有用

前两项主要是给老卡的（新卡上 vLLM 原实现本来就有张量核，默认不替换），第三项哪张卡都能用。在单卡 Blackwell（sm_120，16 GB）上跑 Qwen3.5-4B，只开草稿裁剪这一项，decode 吞吐1.13~1.19×，首字时间不变（说明插件本身没有额外开销）。

精度

• GSM8K 前 40 题：两边都是 39/40，错的是同一道，40 题的最终答案全部相同，33 题输出逐字相同（其余在中途分叉：浮点舍入不同让某一步贪心选了别的词，答案不变）。
• 每个 kernel 都和 fp32 参考结果逐个对比过，误差和 vLLM 原实现在同一量级。
• 每个优化都有生效条件，不满足就自动退回 vLLM 原实现，并在日志里说明原因。

怎么用

环境是 vLLM 0.27.1（官方镜像 vllm/vllm-openai:v0.27.1）。装上 wheel，照常启动：

gitclone https://github.com/ZiMianHe/vllm-boost&&cdvllm-boost pipinstall--no-deps release/v0.1.0/vllm_boost-0.1.0+vllm0.27.1-cp312-cp312-linux_x86_64.whlVLLM_BOOST_EXPERIMENTAL=1vllm serve<模型目录>--tensor-parallel-size2--max-num-seqs8--max-model-len262144\--gpu-memory-utilization0.935--max-num-batched-tokens4096\--enable-prefix-caching --trust-remote-code --language-model-only\--attention-backend TRITON_ATTN\--speculative-config'{"method":"mtp","num_speculative_tokens":3,"attention_backend":"TRITON_ATTN"}'\-cc'{"cudagraph_capture_sizes": [1, 2, 4, 8, 16, 24, 32]}'\--enable-auto-tool-choice --tool-call-parser qwen3_coder --reasoning-parser qwen3 --enable-prompt-tokens-details

关键是主模型和草稿的两处 TRITON_ATTN：用 Triton 后端拿到完整 CUDA Graph，计算由插件换成 FlashInfer。想对比效果，加 VLLM_BOOST=none 就等于没装。

局限

• 目前只支持 vLLM 0.27.1。插件接的是 vLLM 内部实现，换版本要重新适配。
• 主要测了单并发和 4 并发，高并发没有系统测过。
• 端到端只在 2 × 2080 Ti + Qwen3.8-27B 和单卡 sm_120 + Qwen3.5-4B 上验证过，其他卡和模型欢迎反馈。
• MTP接受率在不同场景下差异可能较大。

附录

以下是优化前后实测性能数据。
优化前：

优化后：