大推理模型(LRM)靠生成超长思维链(CoT)提升准确率,代价是推理成本极高(单题平均生成约 1.2 万 token,可到 6.4 万)。NVFP4 是英伟达 Blackwell 上原生支持的 4-bit 微缩放浮点格式,标称峰值吞吐是 BF16 的约 4 倍、权重体积约 1/4,是"低延迟推理服务"最诱人的杠杆——但直接用,两道坎:
NVFP4 属于"微缩放 FP4"格式家族:元素本身只有 4-bit(E2M1),靠共享的缩放因子恢复动态范围。和同家族的 MXFP4 相比,NVFP4 的块更细(16 元素一块)、块缩放用 FP8(E4M3)并额外叠加 FP32 全局缩放——低精度 W4A4 下量级校准更稳,这决定了它是英伟达在 B200 上主推的格式:
| Format | Element Type | Block Scale | Global Scale | Block Size |
|---|---|---|---|---|
| MXFP4 | FP4 (E2M1) | FP8 (E8M0) | – | 32 |
| NVFP4 | FP4 (E2M1) | FP8 (E4M3) | FP32 | 16 |
元素都是 E2M1 的 FP4;MXFP4 块大(32)且块缩放为纯指数格式(E8M0),元数据更紧凑;NVFP4 块小(16)+ 全局 FP32 缩放,精度控制更细,代价是略微多一点的缩放开销。
关于量化精度,论文点了一个重要的反直觉事实:现有 NVFP4 PTQ 方法(旋转、通道缩放、动态块缩放等)全都在压权重量化误差,但 NVFP4 在组大小 16 下权重误差本来就小,这些方法与推理准确率并不干净地相关——与其继续优化误差,不如直接看解码过程本身。这正是 ReSET 的出发点。
推理模型的成本维度是每 token 解码延迟而非聚合吞吐:R1 在 AIME 上平均生成约 12K token(最多 64K),TPOT 直接卡 SLO;再加上 KV-cache 压力、prefill 与 decode 互相干扰、生成停顿,SLO 能容忍的解码 batch 只有 M≤8(B200 上)。而 NVFP4 的 4× 峰值吞吐要靠 Blackwell 的 tcgen05.mma 指令兑现,该指令的 tile 沿 token 维固定为 M=128。下图是这两件事的实证(点击图片可看 300 DPI 原图):
💡 点击任意图片可查看原始高清大图,再次点击或按 Esc 关闭。
作者用 R1-Qwen-14B 在 90 道 AIME(2022–2024)题上生成的 150 万 token 做分析。第一步先把 token 按熵分类:低熵 = 数字、运算符这类"局部没得选"的符号决策;高熵 = "Alternatively" 这类语义灵活的分支点。下图(a)(b)两列就是这两类 token 的典型面孔,而(c)(d)(e)揭示了量化真正做错事的方式:
既然错在低熵符号 token 的偶发误采样,最直接的办法就是对这些位置"降温",把分布压尖。用固定阈值 $\tau_0$(取全体 token 熵的 80 分位,约 0.6)把 token 分成两组:低熵组用低温 $T_{\text{low}}$,高熵组用 $T_{\text{high}}$(公式(1))。结果就是上图 (f) 的曲线:
降 $T_{\text{low}}$(往左看曲线)稳定带来收益,但只能部分恢复准确率;而改 $T_{\text{high}}$ 收益很有限。这说明"低熵 token"这个干预目标找对了,但token 级熵不是正确的控制信号——用固定阈值看单个 token 的熵,会漏判。
为什么会漏判?作者发现粒度错位:现有熵感知解码在 token 级考虑不确定性,但推理中真正波动的不确定性活在 步骤级。所谓"推理步骤",指一段连贯的中间推理单元(实现上用双换行符 \n\n 切分),步骤熵 = 步内 token 熵的平均。下图是全篇最重要的诊断:
再往下挖一层,量化对熵的影响其实是双方向的(附录 B.2 的专门分析):低熵符号 token 的熵被系统性抬高(高不确定步骤里抬得更多)——这就是采样错误变多的直接原因;而极高熵 token($H_t$≥2.0)在高不确定步骤里熵反而下降,典型如 "Alternatively" 这类话语连接词的概率质量被量化压到更少的几个续写上,分支多样性被吃掉。这两个效应合起来,给后面的温度策略提供了完整依据:$T_{\text{low}}$ 压掉"假不确定性",$T_{\text{high}}$ 还回"被压缩的多样性":
ReSET = Reasoning Step Entropy-based Temperature scaling,两个组件:步骤感知阈值(SAT,§4.1)与在线步骤熵估计器(HSE,§4.2)。下面用原文 Figure 4 把整个判定流程走一遍:
固定阈值只在一种情况下失手:高不确定步骤。因此把阈值改为"全局运行均值 $\bar{H}$ 对比 在线步骤熵估计 $\hat{H}_{\text{step}}(t)$"的择一选择:
为什么笃定步骤不能也走"步骤相对"路线?因为在笃定步骤里 $\hat{H}_{\text{step}}$ 退化得极小,相对阈值会反过来误伤正常符号 token——论文在附录里专门用下图说明了这个"退回分支"的必要性:
SAT 在每个解码位置都需要 $\hat{H}_{\text{step}}(t)$,但第 3 节的离线定义要平均"还没生成的 token",不满足自回归的因果性。难点在于步内有两种相反的偏差-方差权衡:步中后段,步内运行平均偏差最低(而且图 4b 显示它稳定优于 EMA——说明"步边界"比"时间近因"信息量大);步开头 w 个位置,样本太少且常是"步开场"话语 token,用最近 w 个 token 的滑动窗口初始化,容忍一点跨步偏差换方差。混合估计器:
图 4b 的横轴正是这几条曲线的对比:窗口初始化(Init)在步前段显著降低方差,步内平均(Causal)在后期偏差最低,EMA 两者都不占优。
| Model | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B |
|---|---|---|---|---|---|
| $\tau$ | 0.6446 | 0.6488 | 0.5505 | 0.4863 | 0.5363 |
$\tau_0$ ≈ 0.5–0.65:Qwen3 系列的熵分位略低,与模型自身输出习惯相关;部署时无需任务级数据即可标定。
每 token 的计算量:一次熵归约 + 两次标量更新($\hat{H}_{\text{step}}$、$\bar{H}$)+ 一次分支——相比完整解码步可忽略,具体开销计入第 6 节的端到端测量。
第 2.2 节已经看到<1% 的利用率,机制在这里:Blackwell 的 NVFP4 GEMM 路径(vLLM / CUTLASS / MR-GPTQ)都走 tcgen05.mma,tile 沿 token 维固定 M=128。看 figure 5a:逻辑激活 X∈ℝM×K 无论 M 是多少都要被填充(pad)成 128×K,只有 M 行输出是真正有用的——M=8 时 6.25% 的 tile 是有效行,其余全是"浪费输出"。吞吐导向的框架靠上千的大 batch 把填充摊薄,但 SLO 可行域里优势归零。
CUDA-core 路径可以在线程级灵活暴露 M,不需要固定 tile——但现有框架并没有 NVFP4 的 CUDA-core 实现,这个空白正是本文要填的。要攻克的难点有三个:(C1) 复用流式加载的权重瓦片——所有 token 共享同一个 W,逐个 token 处理会把同一块权重反复从 HBM 拉取;(C2) 足够的并行线程与独立累加链——朴素映射会让 half2 FMA 沿 K 串成一条依赖链,小 M 下 CUDA core 空转;(C3) FP4 解包 + 共享缩放反量化既不能落地到中间 buffer,也不能在内层 K 循环里插同步。
运行时调度:M=1–2 恒用 CUDA-core 核;M≥128 走自研 Tensor-Core GEMM;中间形状按 (M, shape) 实测延迟二选一。因为 vLLM 按 batch 大小抓 CUDA graph,选择在抓图时固定,回放零分支开销。
| Layer Type | vLLM | MR-GPTQ | Ours | |||
|---|---|---|---|---|---|---|
| Q | GEMM | Q+R | GEMM | Q | GEMV | |
| QKV | 1.81 | 13.41 | 4.65 | 11.84 | 1.65 | 6.62 |
| Out | 2.04 | 12.40 | 4.87 | 11.90 | 1.81 | 7.73 |
| Gate-Up | 1.83 | 25.89 | 4.55 | 25.90 | 1.67 | 13.25 |
| Down | 2.80 | 33.26 | 5.08 | 33.87 | 1.79 | 15.00 |
旋转融合激活量化(Q+R)在 M=1 时每层要 4.55–5.08µs,约为普通激活量化(Q)的 2.3 倍;本文流水线吃非交错缩放布局,激活量化保持在 1.65–1.81µs,投影主体 GEMM→GEMV 后延迟砍半以上。
第一,量化开销是被减少而不是被增加。MR-GPTQ 这类旋转方案的代价是解码步里每层多花 2.3× 的激活量化时间;ReSET 体系里 CUDA-core 核吃非交错缩放布局、ReSET 采样器本身又不做旋转,所以从系统层面直接省掉了这笔开销(Table 3)。第二,M=1 的投影延迟全面优于 vLLM(Table 4,见下节)。设计选择的贡献量在附录 C 里做了消融——C1(CTA 融合)随 M 增大越来越重要,C2(多累加器)在 M=1 也有效:
| Config | M=1 | M=4 | M=8 |
|---|---|---|---|
| Ours: full kernel (C1+C2+C3) | 15.0 | 27.5 | 58.3 |
| without C1 (no CTA fusion) | 15.3 | 35.9 | 70.9 |
| without C2 (single-row threading) | 16.8 | 32.4 | 64.6 |
| without C1 and C2 | 18.7 | 39.4 | 76.5 |
C1 在 M=8 时把 70.9µs 拉回 58.3µs(权重瓦片复用);C2 在 M=1 时也有 1.8µs 的收益(ILP);两者都去掉则 M=8 慢 1.31×。C3 因会改变内核流水结构(共享内存分配、同步位置、流水深度)而不做单独消融。
主结果是下表:三个 benchmark × 五种方法 × 五个模型,ReSET 在每一个 benchmark 的平均分上都排第一,且是唯一在每个任务上都超过 RTN 基线的方法:
| Task | Method | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B | Avg |
|---|---|---|---|---|---|---|---|
| AIME-120 | BF16 Baseline | 45.7 | 57.4 | 70.4 | 76.1 | 75.8 | 65.1 |
| RTN | 39.6 | 52.4 | 62.5 | 70.4 | 74.4 | 59.9 | |
| BRQ | 41.4 | 49.8 | 53.8 | 66.9 | 73.0 | 57.0 | |
| 4/6 | 41.1 | 53.1 | 64.0 | 70.1 | 74.8 | 60.6 | |
| MR-GPTQ | 39.6 | 50.6 | 65.2 | 71.0 | 73.3 | 60.0 | |
| ReSET | 43.8 | 54.0 | 64.9 | 72.1 | 77.5 | 62.5 | |
| GPQA-Diamond | BF16 Baseline | 48.3 | 58.4 | 54.4 | 61.3 | 64.1 | 57.3 |
| RTN | 47.1 | 53.7 | 50.7 | 57.4 | 62.8 | 54.3 | |
| BRQ | 45.5 | 56.7 | 53.5 | 58.2 | 60.8 | 55.0 | |
| 4/6 | 44.6 | 55.4 | 51.9 | 57.6 | 64.6 | 54.8 | |
| MR-GPTQ | 43.4 | 57.4 | 52.4 | 60.9 | 60.6 | 54.9 | |
| ReSET | 46.0 | 57.6 | 53.2 | 58.5 | 62.9 | 55.6 | |
| LiveCodeBench | BF16 Baseline | 28.2 | 38.5 | 43.2 | 46.5 | 49.9 | 41.3 |
| RTN | 29.5 | 37.1 | 36.4 | 46.7 | 46.5 | 39.2 | |
| BRQ | 27.6 | 37.7 | 36.2 | 39.0 | 45.1 | 37.1 | |
| 4/6 | 27.0 | 36.4 | 35.8 | 45.3 | 45.7 | 38.0 | |
| MR-GPTQ | 27.2 | 34.9 | 42.9 | 47.2 | 45.1 | 39.5 | |
| ReSET | 28.4 | 37.9 | 42.1 | 46.1 | 46.7 | 40.2 |
几个值得注意的细节:① 先进 PTQ(BRQ、4/6、MR-GPTQ)在大模型上并不稳定超过朴素 RTN(如 AIME-120 上 MR-GPTQ 只在 Qwen3 系列略好、R1 系列反而不如 RTN),佐证了"NVFP4 下再压权重误差收益甚微"的判断;② ReSET 的最大增益在 AIME-120:AIME 平均 +2.6(Qwen3-32B +3.1、R1-Qwen-7B +4.2),而且 Qwen3-32B 的 77.5 反超了 BF16 基线(75.8)——温度策略不仅补偿量化损失,还顺带改善了采样;③ GPQA / LiveCodeBench 的绝对提升相对温和(平均 +1.3 / +1.0),但仍是全基线最优。
先看阈值机制本身的贡献。在相同解码设置下对比三种阈值策略:固定阈值(公式 1)、滑动窗口阈值、ReSET——差异只在"熵阈值怎么定":
| Method | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B | Avg |
|---|---|---|---|---|---|---|
| BF16 | 45.7 | 57.4 | 70.4 | 76.1 | 75.8 | 65.1 |
| NVFP4 | 39.6 | 52.4 | 62.5 | 70.4 | 74.4 | 59.9 |
| Fixed Threshold | 41.8 | 53.2 | 63.3 | 70.3 | 75.8 | 60.9 |
| Sliding Window | 42.6 | 51.3 | 63.3 | 70.6 | 76.4 | 60.8 |
| ReSET | 43.8 | 54.0 | 64.9 | 72.1 | 77.5 | 62.5 |
固定阈值与滑动窗口对 NVFP4 的提升不稳定(Fixed Threshold 在 Qwen3-14B 上反而 -0.1),ReSET 在每个模型上都是最优——验证"步骤相对熵"才是有效信号。
再看"截断式采样"是否也能替代。有人会问:直接把 top-p 或 min-p 参数扫一遍不就行了?下表是 T=0.6 固定下扫描的结果——top-p 平均只从 59.9 微动到 60.3 封顶,min-p 系列全面低于基线(58.2–59.4):
| Method | T | Top-p | Min-p | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B | Avg |
|---|---|---|---|---|---|---|---|---|---|
| NVFP4 Baseline | 0.6 | 0.95 | 0.0 | 39.6 | 52.4 | 62.5 | 70.4 | 74.4 | 59.9 |
| NVFP4 + Top-p Sweep | 0.6 | 0.80 | 0.0 | 41.0 | 49.5 | 61.0 | 70.1 | 74.3 | 59.2 |
| 0.85 | 39.7 | 50.6 | 63.3 | 70.2 | 74.5 | 59.7 | |||
| 0.90 | 40.2 | 52.4 | 63.0 | 71.2 | 74.5 | 60.3 | |||
| 0.99 | 38.4 | 53.9 | 63.6 | 69.6 | 75.2 | 60.1 | |||
| 1.00 | 40.1 | 52.6 | 63.9 | 70.7 | 73.6 | 60.2 | |||
| NVFP4 + Min-p Sweep | 0.6 | 0.95 | 0.01 | 39.8 | 52.4 | 61.2 | 68.7 | 74.8 | 59.4 |
| 0.03 | 40.3 | 49.7 | 60.6 | 70.0 | 74.5 | 59.0 | |||
| 0.05 | 38.7 | 50.2 | 58.5 | 70.2 | 73.3 | 58.2 | |||
| 0.10 | 41.0 | 51.2 | 57.6 | 69.8 | 73.3 | 58.6 | |||
| ReSET | ReSET | 0.95 | 0.0 | 43.8 | 54.0 | 64.9 | 72.1 | 77.5 | 62.5 |
ReSET 把表 5 的 T=0.6 固定条件放开为自适应温度,全面超过所有截断组合——证明问题在"熵控制"而非"分布截断"。
最后验证"选择性降温"而非"一刀切"。既然低不确定步骤也要管,那对这些步骤里的所有 token 一律用 $T_{\text{low}}$ 会不会更好?答案是否定的(Table 6):
| Policy | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B | Avg |
|---|---|---|---|---|---|---|
| All $T_{\text{low}}$ | 40.6 | 51.2 | 65.1 | 72.0 | 74.1 | 60.6 |
| ReSET | 43.8 | 54.0 | 64.9 | 72.1 | 77.5 | 62.5 |
均匀降温(60.6)比 NVFP4 基线(59.9)只好 0.7 分,而选择性降温(62.5)好 2.6 分:低不确定步骤里并非所有 token 都严格确定,一刀切会压掉合法备选。
内核级。先是 M=1 的逐投影对比(Table 4):12 种投影形状全部快于 vLLM-CUTLASS,1.57–2.49×,最大增益出现在 Qwen3-8B 的 down-projection(N=4096, K=12288)。K 越大的层(如 Gate-Up/Down)加速越明显,因为它们更受"填充浪费 + 权重流式加载"的影响:
| Qwen3 | N | K | vLLM | Ours | Speedup |
|---|---|---|---|---|---|
| 8B | 6144 | 4096 | 8.69 | 4.28 | 2.03× |
| 4096 | 4096 | 7.69 | 3.75 | 2.05× | |
| 12288 | 4096 | 11.45 | 6.40 | 1.79× | |
| 4096 | 12288 | 16.61 | 6.66 | 2.49× | |
| 14B | 7168 | 5120 | 10.79 | 5.53 | 1.95× |
| 5120 | 5120 | 9.08 | 4.64 | 1.96× | |
| 17408 | 5120 | 14.93 | 9.49 | 1.57× | |
| 5120 | 17408 | 23.77 | 10.35 | 2.30× | |
| 32B | 10240 | 5120 | 13.41 | 6.62 | 2.03× |
| 5120 | 8192 | 12.40 | 7.73 | 1.60× | |
| 25600 | 5120 | 25.89 | 13.25 | 1.96× | |
| 5120 | 25600 | 33.26 | 15.00 | 2.22× |
端到端。下面这张图把内核收益放进完整服务流看:B200 上、512 token 输入、输出最长 32K,对比 BF16、vLLM-CUTLASS、vLLM-FlashInfer 与本文方案(Batch=1 和 8):
最后是三个超参的敏感性证据,结论一句话:都不敏感。窗口大小 w 从 16 到 128,平均分只波动 1.1 分(w=32 最优,过大窗口会糊掉步骤局部统计);$T_{\text{low}}$ 在 0.1–0.4 之间只波动 0.4 分;$T_{\text{high}}$ 则是"越大越好、到 1.0 封顶"的一致趋势——这正好呼应熵塌缩分析:量化压缩了多少多样性,就用升高的 $T_{\text{high}}$ 还回去多少:
| w | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B | Avg |
|---|---|---|---|---|---|---|
| 16 | 44.9 | 53.9 | 65.7 | 70.3 | 76.7 | 62.3 |
| 32 | 43.8 | 54.0 | 64.9 | 72.1 | 77.5 | 62.5 |
| 64 | 44.5 | 54.4 | 63.5 | 72.0 | 75.1 | 61.9 |
| 128 | 43.2 | 52.0 | 64.7 | 71.8 | 75.3 | 61.4 |
| $T_{\text{high}}$ | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B |
|---|---|---|---|---|---|
| 0.6 (Baseline) | 39.6 | 50.9 | 61.9 | 71.4 | 74.9 |
| 0.7 | 39.8 | 51.6 | 64.3 | 70.6 | 74.1 |
| 0.8 | 40.4 | 54.8 | 64.5 | 73.7 | 74.9 |
| 0.9 | 43.8 | 54.5 | 64.6 | 71.1 | 75.3 |
| 1.0 (Default in ReSET) | 42.0 | 52.9 | 64.9 | 70.9 | 77.5 |
| $T_{\text{low}}$ | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B | Avg |
|---|---|---|---|---|---|---|
| 0.1 | 42.0 | 52.9 | 64.9 | 70.9 | 77.5 | 61.6 |
| 0.2 | 42.7 | 52.9 | 64.3 | 70.7 | 75.7 | 61.2 |
| 0.3 | 41.4 | 54.0 | 64.1 | 72.1 | 75.6 | 61.4 |
| 0.4 | 43.8 | 52.2 | 64.8 | 71.2 | 75.4 | 61.5 |
三张表合计:默认配置 (w=32, $T_{\text{low}}$=0.1, $T_{\text{high}}$=1.0) 附近性能平坦,部署无需精细调参。
如果 NVFP4 的采样问题如此突出,那为什么常见 PTQ 手段在推理任务上帮不上忙?WikiText-2 困惑度给出一个侧面答案:RTN 与 4/6、BRQ、MR-GPTQ 的困惑度差距在 0.1–1.0 之间,全部远小于 BF16↔RTN 的量化鸿沟(R1-7B 上 25.07→26.23)。NVFP4 下权重误差已经小到"再压也没意义"的地步——推理精度问题的根源在解码采样,不在权重重构:
| Bit-Precision | Method | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B |
|---|---|---|---|---|---|---|
| BF16 | Baseline | 25.07 | 8.91 | 9.72 | 8.65 | 7.61 |
| NVFP4 | RTN | 26.23 | 9.50 | 10.10 | 8.91 | 8.01 |
| 4/6 | 26.19 | 9.49 | 10.08 | 8.86 | 7.98 | |
| BRQ | 27.00 | 9.81 | 10.35 | 9.53 | 8.17 | |
| MR-GPTQ | 27.48 | 9.64 | 10.06 | 8.95 | 7.81 |
有人会想:ReSET 加在 RTN 上就能追上先进 PTQ,那把 ReSET 直接叠到 GPTQ 上会不会更强?结论是不会(Table 12):除了 LiveCodeBench 平均 +1.0 之外,其余基本打平或略降(如 AIME-120:62.5 → 61.6)。作者的猜测:GPTQ 的纠错作用是"压低低熵 token 的非 top-1 概率",而 ReSET 的温度机制已经在做同一件事,两者功能重叠、收益互相吸收——这也再次说明 ReSET 是 PTQ 之上"便宜的最后一公里",不需要再叠昂贵的重建优化:
| Task | Method | R1-Qwen-7B | R1-Qwen-14B | Qwen3-8B | Qwen3-14B | Qwen3-32B | Avg |
|---|---|---|---|---|---|---|---|
| AIME-120 | RTN | 39.6 | 52.4 | 62.5 | 70.4 | 74.4 | 59.9 |
| ReSET | 43.8 | 54.0 | 64.9 | 72.1 | 77.5 | 62.5 | |
| ReSET† | 41.9 | 53.5 | 65.0 | 72.4 | 76.1 | 61.6 | |
| GPQA-Diamond | RTN | 47.1 | 53.7 | 50.7 | 57.4 | 62.8 | 54.3 |
| ReSET | 46.0 | 57.6 | 53.2 | 58.5 | 62.9 | 55.6 | |
| ReSET† | 44.4 | 57.4 | 52.2 | 59.6 | 64.3 | 55.6 | |
| LiveCodeBench | RTN | 29.5 | 37.1 | 36.4 | 46.7 | 46.5 | 39.2 |
| ReSET | 28.4 | 37.9 | 42.1 | 46.1 | 46.7 | 40.2 | |
| ReSET† | 29.5 | 37.3 | 43.8 | 46.7 | 48.8 | 41.2 |
正文只报了 M=1,这里补全另外两个运行形态(附录 C):M=4/8 时 CUDA-core 核仍全面领先(1.19–1.66×,24 个组合全部更快);M=128(Tensor-Core tile 被充分占用)时,自研 GEMM 相对 vLLM-CUTLASS 仍有 1.00–1.34× 的优势。换句话说,整个 M 轴的调度图谱都是正的:
| Model | Layer | N | K | vLLM M=4 | Ours M=4 | Speedup | vLLM M=8 | Ours M=8 | Speedup |
|---|---|---|---|---|---|---|---|---|---|
| Qwen3-8B | qkv_proj | 6144 | 4096 | 8.32 | 6.18 | 1.35× | 8.05 | 6.01 | 1.34× |
| o_proj | 4096 | 4096 | 7.47 | 6.00 | 1.25× | 7.42 | 5.81 | 1.28× | |
| gate/up_proj | 12288 | 4096 | 10.34 | 8.67 | 1.19× | 10.35 | 8.26 | 1.25× | |
| down_proj | 4096 | 12288 | 16.05 | 11.57 | 1.39× | 15.49 | 11.62 | 1.33× | |
| Qwen3-14B | qkv_proj | 7168 | 5120 | 10.92 | 7.40 | 1.48× | 10.09 | 7.42 | 1.36× |
| o_proj | 5120 | 5120 | 9.09 | 6.88 | 1.32× | 8.66 | 6.81 | 1.27× | |
| gate/up_proj | 17408 | 5120 | 14.88 | 11.21 | 1.33× | 13.59 | 11.19 | 1.21× | |
| down_proj | 5120 | 17408 | 23.58 | 15.40 | 1.53× | 22.22 | 15.02 | 1.48× | |
| Qwen3-32B | qkv_proj | 10240 | 5120 | 13.50 | 8.43 | 1.60× | 11.95 | 8.45 | 1.41× |
| o_proj | 5120 | 8192 | 12.31 | 9.31 | 1.32× | 11.92 | 9.46 | 1.26× | |
| gate/up_proj | 25600 | 5120 | 26.08 | 17.51 | 1.49× | 23.14 | 17.33 | 1.34× | |
| down_proj | 5120 | 25600 | 34.14 | 20.54 | 1.66× | 31.94 | 20.56 | 1.55× |
| Model | Layer | N | K | vLLM (µs) | Ours (µs) | Speedup |
|---|---|---|---|---|---|---|
| Qwen3-8B | qkv_proj | 6144 | 4096 | 13.76 | 11.36 | 1.21× |
| o_proj | 4096 | 4096 | 13.79 | 11.23 | 1.23× | |
| gate/up_proj | 12288 | 4096 | 18.85 | 14.78 | 1.27× | |
| down_proj | 4096 | 12288 | 18.91 | 14.82 | 1.28× | |
| Qwen3-14B | qkv_proj | 7168 | 5120 | 13.89 | 11.30 | 1.23× |
| o_proj | 5120 | 5120 | 13.79 | 11.52 | 1.20× | |
| gate/up_proj | 17408 | 5120 | 20.93 | 18.69 | 1.12× | |
| down_proj | 5120 | 17408 | 25.06 | 18.88 | 1.33× | |
| Qwen3-32B | qkv_proj | 10240 | 5120 | 18.91 | 14.94 | 1.27× |
| o_proj | 5120 | 8192 | 16.86 | 12.74 | 1.32× | |
| gate/up_proj | 25600 | 5120 | 25.06 | 24.99 | 1.00× | |
| down_proj | 5120 | 25600 | 33.25 | 24.80 | 1.34× |
M=128 是 Tensor-Core 主场:本文 GEMM 依然不差于 vLLM-CUTLASS,唯一打平的是 Qwen3-32B gate/up(1.00×);整体 1.12–1.34×。
阅读中遇到缩写,可随时回到这里;正文里的缩写也都可以悬停查看释义。
| 缩写 | 全称 | 一句话解释 |
|---|---|---|
| LRM | Large Reasoning Model | 靠生成长思维链解题的大模型,如 DeepSeek-R1、Qwen3 |
| CoT | Chain-of-Thought | 思维链:解题前先生成一步步推理过程 |
| NVFP4 | NVIDIA FP4 (microscaled) | Blackwell 原生 4-bit 格式:FP4 元素 + 组内共享缩放 |
| W4A4 | Weight 4-bit / Activation 4-bit | 权重与激活都量化为 4-bit 的推理模式 |
| PTQ | Post-Training Quantization | 训练后量化:不重训,直接量化权重/激活 |
| RTN | Round-To-Nearest | 最朴素的"就近舍入"量化基线 |
| BRQ / 4-6 / MR-GPTQ | — | 三种先进 PTQ 基线:块旋转 / 自适应块缩放 / 重建式 |
| SLO | Service Level Objective | 服务质量目标,如"单 token 延迟 ≤ 某值" |
| TPOT | Time Per Output Token | 每生成一个输出 token 的耗时,延迟关键指标 |
| KV cache | Key-Value cache | 缓存历史键值对,避免重复计算,吃显存 |
| Tile / M | Tensor-Core tile | Tensor Core 每次矩阵乘的最小块,本文 token 维固定 M=128 |
| GEMV / GEMM | Matrix-Vector / Matrix-Matrix | 向量×矩阵(解码)与矩阵×矩阵(prefill/大 batch) |
| CTA | Cooperative Thread Array | 协作线程数组,即一个线程块 |
| ILP | Instruction-Level Parallelism | 指令级并行:多条独立指令同时执行 |
| Perplexity | 困惑度 | 语言模型对文本的"惊讶程度",越低越好 |