XFP:质量地板驱动的自适应码本量化——3.4 位在双卡工作站跑通 397B MoE

9863 字
49 分钟
XFP:质量地板驱动的自适应码本量化——3.4 位在双卡工作站跑通 397B MoE

AI 生成内容声明

背景:797 GB 的模型,要塞进 192 GB 的工作站#

事情的起点很具体:论文作者 Thomas Witt 想在一台只有两张 RTX PRO 6000 Blackwell(各 96 GB,合计 192 GB)的工作站上跑 Qwen3.5-397B-A17B——一个 397B 参数的混合线性注意力/自注意力 MoE 模型,每层 512 个路由专家、共 60 层,BF16 权重原始体积约 795 GB。795 GB 对 192 GB,差了一个数量级,任何不压缩的方案都不用讨论。

显存装不下,量化是唯一的路。但问题随之而来:INT4 量化后模型约 200 GB,加上 KV cache、激活和运行时缓冲区,192 GB 依然塞不下;就算勉强塞进 122B 级别的模型,4 位量化带来的质量损失也相当可观。作者在论文里把这段经历总结成一句话:真正的难点不是”怎么量化”——GPTQ、AWQ 早就解决了这个问题——而是”为什么 4 位的质量损失这么大”。

这个问题问得很准。如果 4 位量化的损失是可以避免的、是被某种可修正的结构性因素造成的,那么压缩的下限就可以继续下探;如果损失是 4 位本身的信息论极限,那就只能接受。XFP 的答案站在前者一边:损失主要来自”用同一套表示模板去套所有层”,而这一点是可以改的。

在展开 XFP 之前,先把量化领域的三条技术路线和它们各自的隐含假设理清楚,这样才能看出 XFP 到底改了什么。

量化的三条路线,以及它们共享的隐含假设#

路线一:线性量化(GPTQ / AWQ / AutoRound)#

线性量化的核心是把权重映射到一个均匀网格上。一个权重 ww 的量化与反量化可以写成:

w^=sround(wzs)+z\hat{w} = s \cdot \text{round}\left(\frac{w - z}{s}\right) + z

其中 ss 是量化步长(scale),zz 是零点(zero point),w^\hat{w} 是重建值。sszz 通常按组(group)或按通道(channel)计算,用局部的最大值或分位数来吸收不同位置权重的幅值差异——这就是 per-group / per-channel scaling。整数档位有 2N2^N 个(4 位就是 16 档),一旦 ss 定死,网格的分辨率就定死了。

这条路线上的三个代表作各用不同方式补偿量化误差:

  • GPTQ 用 Hessian 矩阵(二阶导信息)逐列做误差补偿:量化完一列权重后,把这一列引入的误差按 Hessian 逆近似分摊到尚未量化的列上,让后续列”吸收”前面列的错误。代价是需要校准数据和逐层的矩阵求逆计算。
  • AWQ 观察到不同权重通道对激活的影响不同(salient channels),于是按激活统计量给重要通道乘上保护性缩放,把量化误差从重要通道转移到不重要的通道。它不需要 Hessian,但仍需要校准数据。
  • AutoRound 把”四舍五入”改成”学习舍入”:对每个权重学习”向上还是向下取整”的二元决策,用带符号梯度下降在小校准集上优化,几分钟内就能得到接近无损的 INT4 质量。独立评测显示 AutoRound INT4 的精度经常持平甚至超过有专用硬件加持的 NVFP4。

AWQ 的激活感知显著通道保护
AWQ 的激活感知显著通道保护

图:AWQ 的激活感知量化思路——只保护约 1% 的显著权重通道(salient channels):不是保留 FP16,而是给这些通道乘保护性缩放、同时反缩放对应激活,把量化误差转移到不重要的通道。(图源:AWQ 论文 / llm-awq 官方仓库 overview 图)

三种方法共同的隐含假设是:所有层、所有专家、所有投影矩阵都使用同一个均匀网格模板。离群值(outlier)不是被分离出来,而是靠调整 ss 或误差补偿去”消化”。这在离群值温和时成立,在离群值极端时失效——下面用具体数字说明为什么。

设想一个权重矩阵,99% 的权重落在 ±1σ\pm 1\sigma 之内,但存在一个 49σ49\sigma 的极端离群值(这不是虚构的数字:论文在 GLM-4.7-Flash 的 kv_a_proj 层实测到 49σ49\sigma 的极值,见论文附录 B)。为了让这个离群值能被表示,ss 必须覆盖 ±50σ\pm 50\sigma 的范围。4 位有 15 个非零档位,步长约为 100σ/156.7σ100\sigma / 15 \approx 6.7\sigma——也就是说,落在 ±1σ\pm 1\sigma 区间内的绝大多数权重,离它们最近的档位只有”零”这一个档位。一个离群值把整个矩阵的主体权重压成了零。per-group scaling 能缓解但不能根治:离群值所在的组(128 个权重)整组遭殃。论文把这称为”离群值驱动的缩放让主体分布只剩一个可用的非零档位”,这是 4 位质量损失的第一大来源。

路线二:硬件格式(NVFP4 / MXFP4)#

第二条路线是给硬件定制的对数式格式。NVIDIA 在 Blackwell 上主推的 NVFP4 把 4 位拆成 1 位符号 + 2 位指数 + 1 位尾数(E2M1),每 16 个权重共享一个 FP8 缩放因子。E2M1 的所有取值组合恰好构成一个固定的 16 元码本——所以 XFP 论文说”NVFP4 的浮点之名指的是码本条目,而不是表示本身”。MXFP4 的思路类似:每个元素 1 位符号 + 2 位尾数(E2M1 形式),每 32 个元素共享一个 8 位指数,由 OCP 的 Microscaling 规范定义。

NVFP4 与 MXFP4 的块缩放精度对比
NVFP4 与 MXFP4 的块缩放精度对比

图:NVFP4 与 MXFP4 的块缩放对比——NVFP4 用 FP8 E4M3 分数缩放因子(可贴合每个 block 的取值分布),MXFP4 用 2 的幂缩放(UE8M0),分数缩放的量化误差更低。(图源:NVIDIA 官方博客 Introducing NVFP4)

这类格式的问题有两个,而且都很要命:

第一,码本是全宇宙统一的。 同一个 16 元码本套在每一个层、每一个模型、每一种架构上,不管这一层的权重分布是什么形状。注意力层那 0.3% 的离群值吃掉大部分表示预算,高斯分布的专家层却被过度供给——问题与路线一一模一样。

第二,硬件加速只存在于数据中心。 NVFP4 的专用 Tensor Core 路径只在数据中心级 Blackwell(SM100+,即 B100/B200)上可用。XFP 论文给出了工作站级 Blackwell(SM120/SM121)上的实际生态状况,信息量很大:CUTLASS 的 NVFP4 MoE kernel 在 SM120 上因 TMA work-sharing 的 grouped-GEMM 初始化错误直接失败;FlashInfer 无法为 sm120f 家族目标编译(FP4 转换 PTX 指令所需的 target 不支持);vLLM 的后端选择逻辑不认识 SM120 是”NVFP4 可用”的架构,静默回退到 Marlin INT4。社区陆续提交补丁之后,NVFP4 checkpoint 总算能在 RTX PRO 6000 和 DGX Spark 上跑起来,但即使在修复之后,SM120 上的原生 NVFP4 Tensor Core 路径在 MoE 负载上仍然跑不过 Marlin W4A16 回退路径,因为 grouped-GEMM 的特化尚未完成。

结论很反直觉但很清晰:在最需要压缩的工作站层级,NVFP4 的硬件优势恰好不可用;而它那个固定码本在任何硬件上都一样粗糙。独立评测(如 The Kaitchup 的 NVFP4 实测)也普遍发现 NVFP4 的精度不如同位的校准感知方法(AutoRound、AWQ)——它的优势是吞吐,而这个吞吐优势在工作站上拿不到。

路线三:码本量化(AQLM / QuIP# / SqueezeLLM / K-quants)#

第三条路线放弃均匀网格,改用学习出来的码本:每个权重不存数值,只存一个指向码本条目的小整数索引。这类方法在学术上做到了 2 位的极限压缩,但各有代价:

  • AQLM 把权重向量表示为多个(通常 2 个)学习码本条目的和,2 位质量很好,但多码本优化依赖校准数据且计算昂贵。
  • QuIP#/QTIP 先用随机正交变换(incoherence processing)把离群值的能量”摊平”到所有权重上,再做 2 位量化;代价是 Hessian 计算和解码端的正交变换开销。
  • SqueezeLLM 和 SpQR 与 XFP 在结构上最接近:都是”稀疏高精度分量 + 密集低精度分量”的分解。区别在于判定准则——SpQR 用 Hessian 判断哪些权重”敏感”,XFP 直接用权重分布本身(离群值标准)判断,把 Hessian 完全消掉。另一个区别是密集分量的形式:SpQR 用固定线性量化器,XFP 用每组学习码本。
  • 工程上最普及的码本量化其实是 llama.cpp 的 K-quants(Q2_K 到 Q6_K):分块码本、CPU 推理,是”分块码本量化在大规模部署中”最接近 XFP 的先例,但 K-quants 的码本是固定公式生成而非学习的,也没有离群值分离。

SqueezeLLM 的 Dense-and-Sparse 分解
SqueezeLLM 的 Dense-and-Sparse 分解

图:SqueezeLLM 的 Dense-and-Sparse 分解——权重矩阵被拆成可激进量化的密集码本分量 + 保留敏感/离群权重的稀疏 FP16 分量,与 XFP 的 Woutlier+Wcodebook\mathbf{W}_{\text{outlier}} + \mathbf{W}_{\text{codebook}} 结构同源。(图源:SqueezeLLM 论文 / 官方仓库)

三个结构性问题#

把三条路线的问题归纳起来,XFP 论文给出了三个结构性问题(P1-P3),整个方法就是围绕它们设计的:

  • P1 离群值问题:Transformer 权重矩阵是重尾分布,0.3%–2% 的权重占据远超主体的幅值范围。线性量化靠单个 scale 容纳整个范围,牺牲主体分辨率;极端情况下主体只剩一个可用非零档位。
  • P2 MoE 异质性问题:MoE 模型里不同组件类别的权重分布差异巨大。实测发现路由专家(routed experts)对粗量化的容忍度远高于注意力和共享专家——一个全局统一的质量标准无法利用这种不对称性。
  • P3 码本存储问题:每通道学习码本占用的共享内存(SMEM)与输出维度成正比。工作站 GPU 每个 CTA 只有 99 KB SMEM,这会挤占激活行缓存,迫使 kernel 过早拆分 KK 维,限制单流解码可达的模型规模。

核心思想:把质量当约束,位宽是算出来的#

XFP 的核心设计决策只有一句话,但这句话把整个工作流倒了过来:

操作者不再指定位宽,而是指定重建质量的下限;位宽变成每层自动推导的结果。

具体来说,操作者给出两个”质量地板”(quality floor)——注意力与共享专家投影的严格地板 τstrict\tau_{\text{strict}},以及路由专家的宽松地板 τlazy\tau_{\text{lazy}}。质量用什么度量?逐通道余弦相似度(per-channel cosine similarity)。XFP 对每一层自动搜索满足地板的最小码本尺寸 NN,把达不到质量要求的权重分离进稀疏 FP16 离群值残差,然后按 NN 位打包。流程可以概括为:质量地板(输入)→ 码本学习 → 余弦检验 → 位宽(输出)→ 打包 → 融合 kernel。

XFP 管线概览
XFP 管线概览

图:XFP 管线概览——操作者只指定质量地板 τ,离群值提取、Lloyd 码本学习、自动选位与融合解码均由算法自动完成。(图源:XFP 论文 Figure 1)

这个反转带来三个直接后果:位宽成为逐层派生量;离群值残差不再是启发式补丁,而是结构上保证质量地板成立的组成部分;同一个算法不加任何人工调参就能适配不同模型——唯一需要操作者调的旋钮是阈值间隔(threshold gap),用来在固定部署目标上做”内存换质量”的权衡。

XFP 的分解形式写作:

WWoutlier+Wcodebook\mathbf{W} \approx \mathbf{W}_{\text{outlier}} + \mathbf{W}_{\text{codebook}}

其中 Woutlier\mathbf{W}_{\text{outlier}} 是稀疏 FP16 矩阵,保存高幅值权重及其原始位置;Wcodebook\mathbf{W}_{\text{codebook}} 是密集的 NN 位索引矩阵,索引一个学习出来的码本。与 SpQR 的关键区别再强调一次:SpQR 的稀疏分量由 Hessian 敏感性决定,XFP 的稀疏分量由分布本身(kσk\sigma 标准)决定——Hessian 被完全消掉。

一个定位表能最清楚地看出 XFP 在方法谱系里的位置:

特性GPTQAQLMQuIP#SpQRNVFP4XFP
学习码本
无 Hessian
无校准数据
离群值分离
自动选位宽
硬件无关
MoE 原生

XFP 用码本学习换取的是”给定质量地板下的最小位宽”,而不是”给定位宽下的最优重建”——这是它和 GPTQ/AQLM/QuIP# 的优化目标差异;用软件定义的吞吐换取硬件无关性,这是它和 NVFP4 的差异。它明确不追求单点最优,追求的是自动、无校准、跨模型一致的工作流。

格式详解:V2 每通道码本与 V2a 共享码本库#

XFP 引擎支持两种码本存储模式,共享同一个自动选位前端和同一个融合解码 kernel 骨架。理解这两种模式是理解全部实验数据的前提。

V2:每通道 Lloyd 码本#

V2 模式下,每个输出通道拥有一个独立的 FP16 码本,包含 2N2^N 个条目,通过逐通道 Lloyd 迭代学习(NN 支持 2 到 6)。每层的存储构成三部分:

  • 打包索引:每个权重 NN 位;
  • 码本:每输出通道 2N2^N 个 FP16 条目;
  • 离群值三元组(行号 int64, 列号 int64, 值 fp16),每个离群值 18 字节。

在默认的 2% 离群值上限下,离群值存储为每权重增加 0.02×18=0.360.02 \times 18 = 0.36 字节。这个数字后面还会出现——它是”位宽 vs 离群值”权衡的定价基准。

子字节打包#

NN 位索引不是逐元素存满字节,而是按机器字打包。打包方案如下表:

NN每字容纳值数存储字使用位保留位V2V2a
216uint32320
310uint32302
48uint32320
53uint16151
65uint32302

V2a 有一个 V2 没有的硬约束:128 权重组必须能被”每字值数”整除。默认组大小 g=128g=128 下,N=3N=3 因为 128mod100128 \bmod 10 \neq 0 被排除;N=5,6N=5,6 则超出 warp 级查表(LUT)预算。所以 V2a 的候选集只剩 {2,4}\{2,4\}——这个看似无伤大雅的细节,正是后面”双阈值”设计诞生的直接原因。

V2a:共享码本库#

V2a 把”每通道一个码本”换成”每层一个共享码本库”:L=32L=32 个码本(默认值),每个 128 权重组分配一个 5 位索引指向最匹配的库条目,外加每组的 scale 和中点(mid-point)。

动机是纯工程性的——SMEM 预算。工作站级 Blackwell(SM120/SM121)每 CTA 只有 99 KB 可选 SMEM,不到数据中心 SM(228 KB)的一半。V2 的每通道码本占用的 SMEM 与输出维度成正比;而 V2a 一个 L=32L=32N=4N=4 的码本库只有 32×16×2=132 \times 16 \times 2 = 1 KB,与维度无关。省下的预算把激活行缓存的上限从 V2 的 K8192K \le 8192 抬到 K32768K \le 32768,代价是每组的码本精度略降(一组共享一个库条目,而不是每通道专属码本)。

码本库的学习方式(LibFit)也值得一提:先按 V2 的方式学出每通道 Lloyd 码本,然后对码本本身再跑一轮 Lloyd 迭代,把全层的码本浓缩成 L=32L=32 个代表性码本。这是”码本的码本”——两级聚类的思想。

Kernel 侧的 lane 几何遵循同一套算术:LANES_PER_GROUP=128/每字值数\text{LANES\_PER\_GROUP} = 128 / \text{每字值数}N=4N=4 时每组 16 条 lane、每次 warp 迭代处理 2 个组;N=2N=2 时每组 8 条 lane、每次处理 4 个组。

有效位宽的账可以精确算出来:索引 NN 位,加上每 128 权重分摊的 5 位库分配 + 16 位 scale + 16 位中点:

beff=N+5+16+16128N+0.29b_{\text{eff}} = N + \frac{5 + 16 + 16}{128} \approx N + 0.29

论文给出的数字是 V2a 在 N=2N=2 时约 2.31 有效位、N=4N=4 时约 4.31 有效位(均不含离群值残差)。这个 0.29 位的常数开销,就是”共享库”模式的固定入场费。

编码管线:离群值分离、Lloyd 迭代与自动选位#

离群值提取#

对每个权重矩阵 W\mathbf{W}:先计算均值 μ\mu 和标准差 σ\sigma,把满足以下条件的权重提取进稀疏 FP16 残差:

wμ>kσ,k=4.0|w - \mu| > k\sigma, \quad k = 4.0

默认 k=4.0k=4.0,且提取比例硬性封顶 2%(超过时只取幅值最大的 2%)。被提取的位置在主体矩阵中μ\mu 回填——这个细节很重要:回填均值而不是零,能让主体分布保持完整形状,Lloyd 迭代面对的是一份”干净”的数据。k=4.0k=4.0 的选择也符合直觉:高斯分布下 z>4σ|z|>4\sigma 的概率只有约 0.006%,所以对接近高斯的层几乎不触发;只有重尾层才会真正用到 2% 的上限。

Lloyd 迭代#

码本学习用的是教科书级的 Lloyd 算法(Lloyd-Max 量化器,1982 年 Lloyd 的经典论文,本质是标量版 k-means)。初始化不是随机采样,而是按 CDF 均匀分位点2N2^N 个初始条目——每个初始条目拥有相等的概率质量,这比随机初始化稳定得多。随后交替执行 20 轮”最近邻分配 + 条件均值更新”,收敛到局部最优码本:

codebook=argminCimincC(wic)2\text{codebook} = \arg\min_{C} \sum_{i} \min_{c \in C} (w_i - c)^2

迭代轮数有讲究:V2 默认 5 轮,但 V2a 的 MoE LibFit 用 20 轮(XFP_MOE_LLOYD_ITERS=20)。原因很具体:欠收敛的质心会系统性低估 xfp2 的质量,把自动选位推向更保守的 xfp4——20 轮是 397B 模型上让选位稳定的经验值。

自动选位(Auto-Select)#

自动选位算法(论文 Algorithm 1)把质量地板变成位宽决策:

def auto_select(W, layer_class, candidates, tau_strict, tau_lazy):
W_bulk, W_outlier = extract_outliers(W, k=4.0, cap=0.02) # 分离离群值
tau = tau_lazy if layer_class in ROUTED_EXPERT else tau_strict
for N in sorted(candidates): # V2: {2,3,4}; V2a: {2,4}
cb = learn_codebook(W_bulk, 2**N) # Lloyd 迭代(V2a 走 LibFit)
W_hat = reconstruct(cb, W_outlier) # 码本查表 + 离群值回填
cos = per_channel_cosine(W, W_hat)
if median(cos) >= tau: # 中位数门控
return N
return max(candidates) # 兜底:全候选都不达标

逐通道余弦相似度定义为(对每个输出通道 ii):

cosi=Wi,:W^i,:Wi,:2W^i,:2\cos_i = \frac{\mathbf{W}_{i,:} \cdot \hat{\mathbf{W}}_{i,:}}{\lVert \mathbf{W}_{i,:} \rVert_2 \, \lVert \hat{\mathbf{W}}_{i,:} \rVert_2}

两个设计细节值得解释”为什么”:

  • 为什么用中位数聚合而不是均值或最小值? 均值对离群通道敏感——个别重建极差的通道会拉低全层分数;最小值则过于保守,在 N=2N=2 时几乎所有层都过不了关。中位数是一个稳健的折中:一半通道达标即算通过。
  • 为什么是余弦相似度而不是 MSE? MSE 对幅值敏感,余弦对方向敏感。在 GEMM y=xW\mathbf{y} = \mathbf{x} \cdot \mathbf{W} 中,方向保真比幅值保真更关键——激活侧有可学习的 norm 吸收幅值缩放。论文给了一个很有说服力的经验证据:路由专家上 XFP2 的 MSE 比 XFP4 高 12 倍,但数学题的准确率完全相同。

双阈值:一个打包约束逼出来的设计#

两阈值机制不是拍脑袋的设计,而是被 V2a 的打包约束逼出来的。早期 XFP 用单一阈值(τ=0.98\tau = 0.98),在 35B 和 122B 上完全够用——因为 V2 的候选集 {2,3,4}\{2,3,4\} 里有 N=3N=3 这个中间选项:路由专家落到 xfp3、注意力升到 xfp4,没有哪个类别被迫在”过粗”和”过精”之间二选一。

但 V2a 拿掉了 N=3N=3(128 权重组的打包约束),候选集塌缩为 {2,4}\{2,4\}。在 397B 模型上这意味着每个层类别面对二元选择:约 2.31 有效位或约 4.31 有效位。单阈值下只有两种结局——100% xfp4(192 GB 显存放不下)或接近 100% xfp2(数学能力崩溃)。双阈值把这个二元困境解耦成两个统治域:

  • τstrict\tau_{\text{strict}}(默认 0.96)统治注意力、线性注意力和共享专家投影;
  • τlazy\tau_{\text{lazy}}(默认 0.93)统治路由专家路径。

阈值间隔控制路由专家中 xfp2 与 xfp4 的比例,是操作者在固定部署目标上调节”内存 vs 质量”的唯一旋钮。环境变量为 XFP_MIN_COS_STRICT / XFP_MIN_COS_LAZY,机制对 V2 和 V2a 通用。

MoE 采样:4 个专家代替全量#

397B 每层 512 个专家,对每个专家跑 Lloyd + 余弦检验的开销不可接受。XFP 的做法是只对 4 个拼接起来的专家跑自动选位,把选出的位宽应用到全部专家。这个近似在三个模型上做了离线验证:GLM-4.7-Flash(64 专家)0% 不一致、Qwen-35B(256 专家)0% 不一致、Qwen-122B(256 专家)11.5% 不一致——而不一致的位宽决策对数学能力零影响,因为稀疏离群值路径吸收了误差。178 个已验证的 MoE 块中,在线位宽决策与离线预测 100% 一致。

编码成本#

自动选位在 1 分钟内完成。完整的一次量化加载(Lloyd 拟合/LibFit + 离群值提取 + 子字节打包):Qwen3.5-35B 约 10 分钟、122B 约 25 分钟(V2);122B 约 30 分钟、397B 约 55 分钟(V2a),单张 RTX PRO 6000 上完成。打包结果缓存到磁盘,之后加载是 I/O 瓶颈。作为对照,GPTQ 的校准时间与之相当或更长,还多出数据依赖。对”量化加载”这种每次部署都要跑一遍的流程,分钟级且无校准数据的特性是实打实的运维优势。

融合解码 Kernel:一次读取,四步融合#

XFP 的解码 kernel 把每层前向的四个操作融合成一个 pass:

  1. depack:把 NN 位打包索引解包成 int8;
  2. gather:从共享内存中的码本按索引查表,得到 FP16 权重;
  3. scatter-add:把稀疏离群值按位置加回权重矩阵(fp16 输入转 bf16);
  4. GEMM:用 bf16 权重与激活做 Tensor Core 矩阵乘。

码本每个 thread block 只加载一次到 SMEM,同一套 gather 路径服务 N=2N=2N=6N=6 的所有位宽——kernel 是位宽可移植的,这也是两种存储模式能共享一个 kernel 骨架的原因。

// XFP 融合解码 kernel 的骨架(概念示意,非论文原文)
// 每 CTA:1) 载入激活行 x[K] 与码本到 SMEM
// 2) 对每个输出通道组:depack 索引 -> SMEM gather 权重
// 3) 离群值 scatter-add(从全局内存读取 18B 三元组)
// 4) bf16 GEMM(Tensor Core),一次完成
__global__ void xfp_decode_kernel(const PackedIdx* idx, const fp16* codebook,
const OutlierTriple* outliers,
const bf16* x, bf16* y) {
// SMEM: act_row[K_bf16] + cb_tile[chunk * 2^N] + idx_staging
__shared__ bf16 act_row[MAX_K];
__shared__ fp16 cb_tile[CHUNK * (1 << N)];
__shared__ int8 idx_staging[CHUNK * GROUP];
// 载入激活行一次;迭代输出组:depack -> gather -> scatter -> GEMM
// M=1 时权重字节只读一次,激活行全程驻留 SMEM
}

两个工程决策各有血泪教训:

  • bf16 原生路径。早期原型在 fp16 和 bf16 之间来回转换,转换开销让每次 kernel 调用多出 124% 的时间。最终版本激活输入输出全程 bf16,权重 gather 出 fp16 后一次性转 bf16,彻底消掉往返转换。
  • 激活行缓存。单流解码(M=1M=1)是自回归的主战场。kernel 把完整的 KK 维激活行在 SMEM 里缓存一次,然后所有输出组的码本查表都对着缓存的行做点积——激活不再重复从全局内存读取,每个权重字节在解码时恰好被读一次。

M=1 时的内存流量分析是理解一切吞吐数字的钥匙。解码阶段每生成一个 token,模型只激活约 10B 参数(Qwen3.5-122B-A10B 的 A10B 含义),每个权重读一次,KV cache 之外的主要流量就是权重字节。在内存带宽受限的 M=1M=1 场景下,每 token 读多少权重字节,直接决定 tok/s——XFP 自动选位在 122B 上给 GatedDeltaNet 和共享层分配 xfp3(3 位)、路由专家分配 xfp4(4 位),而 Marlin INT4 对所有层统一读 4 位索引 + 每 128 权重一个 fp16 缩放(折合 4.125 位/权重)。解码关键路径(密集层)上 3 位对 4.125 位,字节数少约 27%。这是 XFP 吞吐优势的结构性来源。

K \le 8192 的限制也出自这个设计:V2 的每通道码本随输出维度增长,激活行(K=8192K=8192 时约 24 KB)驻留后 SMEM 即告罄。K>8192K > 8192 的层(397B 的线性注意力投影就是)走 split-KK 变体,把激活行切分到多个 CTA,代价是单流延迟增加 10–20%。V2a 的 1 KB 共享库把上限抬到 K32768K \le 32768,避开 split-KK。V2a kernel 实际以三个 dispatched 变体发布:split-NNM=1M=1 解码)、split-MMM16M \ge 16 批量 prefill)、split-KKK>8192K>8192)。

kernel 以 torch.library.custom_op 封装,兼容 torch.compile 与 CUDA Graph 捕获,并通过 MultiQuant 引擎接入 vLLM——一个按组件类别(注意力、路由专家、共享专家、LM head)注册分发的量化引擎,XFP 的子 4 位操作正是靠这种逐组件策略才成立:在 3 位或 2 位下,“能用”和”崩溃”的边界由最差的组件类别决定。

实验一(Front A):122B 级的质量验证与吞吐#

平台与协议#

两个工作站级 Blackwell 平台:DGX Spark(GB10、SM121、统一内存,单 GPU)与 RTX PRO 6000 Blackwell(SM120、2×962\times96 GB、TP=2)。模型覆盖 GLM-4.7-Flash(30B、MoE-64、纯自注意力)、Qwen3.5-35B-A3B(35B、MoE-256)、Qwen3.5-122B-A10B(122B、MoE-256、GatedDeltaNet + 自注意力混合)走 V2;Qwen3.5-397B-A17B(397B、MoE-512)走 V2a。精度评测用 lm-eval-harness 0.4.11:GSM8K 5-shot、1319 题/seed、3 个 seed,同时报告 flexible-extract 和 strict-match。

一个容易被忽略的合规性检查:所有 XFP 配置都用 FP8 KV cache 和 FP8 LM head。论文先验证了这条基础设施不引入混淆损失——Qwen3.5-35B 在 bf16 与 fp8 KV/LM-head 管线下的 GSM8K strict-match 分别为 76.12% 和 76.35%(Δ=0.2\Delta = 0.2pp,在单 seed 方差内),FP8 基础设施是精度中性的。

主结果#

配置有效位tok/sGSM8K strict
Qwen3.5-122B-A10B,TP=2≈3.9713894.49 ± 0.57
Qwen3.5-122B-A10B,TP=1≈3.9711094.62 ± 0.67
Qwen3.5-35B-A3B,TP=1≈4.03777.18 ± 0.72

122B 模型在约 3.97 有效位下跑到 138 tok/s,GSM8K 94.49% 的 strict-match 与 BF16 基线同级——质量验证通过,压缩率约 4 倍(122B × 2 字节 ≈ 244 GB → 约 60 GB,TP=2 下每卡 57/96 GiB)。

自动选位有个容易被忽略的性质:它在不同模型上收敛到不同的、各自合理的剖面:GLM-4.7-Flash 上注意力 233 层选 xfp3、2 层升 xfp4,路由专家和共享专家全 xfp3,整体约 3.0 有效位;Qwen3.5-122B 上 GatedDeltaNet 和共享层选 xfp3、路由专家升 xfp4(其分布更宽),整体约 3.97 有效位。算法完全一致,结论不同——因为分布不同。这正是”按分布说话”的量化器应有的行为。

离群值分离的消融#

离群值分离到底起了多大作用?论文在 GLM-4.7-Flash 上做了逐层类别的消融(XFP3、k=4k=4),报告”bulk 码本重建”与”加离群值路径后”的余弦与 MSE 变化:

层类型层数cos(仅码本)ΔcosMSE 比(p50)
attn_kva480.981+0.0091.90×
attn_kvb480.983+0.0041.24×
shared_gate_up940.981+0.0021.13×
routed_gate_up960.982+0.00031.02×
routed_down960.982+0.00031.02×

MSE 比 = 仅码本重建的 MSE / 加离群值路径后的 MSE(>1× 表示离群值分离降低了误差)。结论一目了然:离群值分离对注意力层是决定性的(MSE 改善 1.90×),对路由专家几乎无影响(1.02×)。跨模型对比更说明问题:GLM 把灾难性离群值(49σ49\sigma)集中在 kv_a_proj 少数层,Qwen 把温和的重尾均匀分布在各层;相应地,Qwen 的路由专家对离群值分离的响应强度是 GLM 的 6.7 倍。“算法一致、剖面不同”的根源就是分布不同。

对 Marlin INT4:+49% 的受控对比#

吞吐的基准是 Marlin——NVIDIA GPU 上优化得最充分的 INT4 反量化 kernel,服务 AutoRound INT4 checkpoint。同一台 RTX PRO 6000、TP=1 的受控测量:

格式TPtok/sGSM8K strict位宽
XFP V2(auto)111094.62 ± 0.67≈3.97
XFP V2(auto)213894.49 ± 0.57≈3.97
Marlin INT4(AutoRound)173.695.27 ± 0.124.0

TP=1 单流解码 XFP 比 Marlin 快 49%(110 vs 73.6 tok/s),精度差 0.65pp,落在 Lloyd 初始化抖动的包络内(同规模下约 0.67pp 标准差);TP=2 扩展到 138 tok/s,相对 TP=1 Marlin 是 +87%。DGX Spark 上的跨平台数据给出同样的方向:长/中输出上 XFP 比 Marlin 快 16–18%(29.9 vs 25.8、34.8 vs 29.6 tok/s)。V2a 存储模式的探针测量中,122B 上 108.1 tok/s 对 Marlin 73.6(+47%),GLM-4.7-Flash 上则基本打平(118.1 vs 121.1,−2.5%)。

XFP 对 Marlin INT4 的受控对比
XFP 对 Marlin INT4 的受控对比

图:Qwen3.5-122B-A10B 上 XFP 对 Marlin INT4 的受控对比——柱状为单流解码 tok/s,圆点为 GSM8K strict-match(3 seeds,均值 ± std)。(图源:XFP 论文 Figure 3)

有两个诚实性说明要放在这里。其一,论文把 +49% 归因于”3.97 vs 4.0 有效位 → 每 token 读更少字节”,但这个解释并不完整:全模型 3.97 对 4.0 只差 0.75% 的字节,单靠它解释不了 49%。更合理的拆解是分路径看——解码关键路径(GatedDeltaNet + 注意力 + 共享层)上 XFP 是 3 位索引,Marlin 是 4 位 + 缩放因子(4.125 位/权重),字节差约 27%;再叠加融合 kernel 免去 Marlin 的反量化-重打包流水和 bf16/fp16 往返转换。论文没有逐项拆解这 49% 的构成,读者应把该数字理解为”这套 kernel 实现在这套硬件上的实测差距”而非纯位宽收益。其二,Marlin 侧的 GSM8K 数字来自 AutoRound 校准配置,论文明确声明这是一份记录性数据点,不是对两种方法的裁决。

另一个有意思的观察:138 tok/s 时 GPU 利用率 98% 但功耗只有 240 W(TDP 600 W)。高利用率 + 低功耗是内存带宽受限的典型指纹——SM 在等 DRAM,Tensor Core 大量闲置;单流路径里 kernel 启动开销主导,计算侧仍有大量未被利用的余量。这也解释了为什么论文反复声明”大规模并发批量服务不在本文范围”:并发场景需要完全不同的 kernel 调度策略。

实验二(Front B):H-Process——在 397B 上做三约束搜索#

Qwen3.5-397B-A17B 是块硬骨头:512 路由专家 × 60 层,BF16 约 795 GB。任何常规 4 位格式(约 200 GB + KV cache + 激活 + 缓冲)都进不了 2×962\times96 GB。Front B 展示的是 H-Process:对两个余弦阈值做质量驱动的迭代搜索,找到”模型刚好塞进去、输出仍然像话”的工作点。

三约束定义搜索空间#

搜索空间由三条边界界定,缺一不可:

  1. 操作者阈值对 (τstrict,τlazy)(\tau_{\text{strict}}, \tau_{\text{lazy}})——位宽由它们派生,不是直接指定;
  2. OOM 边界(加载期):阈值过高 → 更多路由专家组升入 xfp4 → 编码期的 Lloyd 迭代内存尖峰冲破 96 GB/卡。关键在于尖峰发生在量化加载阶段而非服务稳态——把 gpu_memory_utilization 调到 0.97、KV cache 压到 2 GB、max_model_len 砍到 8192 都救不了,因为服务侧分配还没开始就已 OOM。H1.7(τlazy=0.935\tau_{\text{lazy}}=0.935)三次加载尝试全部死在编码尖峰上;
  3. 垃圾输出边界(生成期):阈值过低 → 路由专家接近全 xfp2 → 位宽重建在数值上通过余弦门,但生成退化成重复 token。余弦相似度是指示器,不是质量保证——垃圾边界严格位于”余弦通过区”的内部,是”包络必须靠评测验证,不能只靠指示器”的经验提醒。

服务配置本身也值得记录:TP=2、BF16 KV cache(FlashInfer 的 FP8-KV 与该模型 head_dim=256 的投影不兼容,Triton 注意力后端 + BF16 KV 是能跑通的组合)、XFP_GROUP_SIZE=128XFP_MOE_LLOYD_ITERS=20、LM head FP8、max_model_len=65536、关闭 custom all-reduce(SM120 在 TP=2 下的已知问题)。

Terminal window
# H-Process 搜索中的关键环境变量(论文 §6.3)
XFP_MIN_COS_STRICT=0.96 # 注意力/线性注意力/共享专家:严格地板
XFP_MIN_COS_LAZY=0.93 # 路由专家:宽松地板(H1.5 工作点)
XFP_GROUP_SIZE=128
XFP_MOE_LLOYD_ITERS=20 # V2a MoE LibFit 的 Lloyd 迭代轮数

质量梯度:窄窗口里的精细调节#

τstrict\tau_{\text{strict}} 全程钉在 0.96,扫 τlazy\tau_{\text{lazy}}

配置τlazy\tau_{\text{lazy}}稳态显存/卡加载尖峰有效位探针(n=50)全量 1319
G0.8070 GiB96 GiB≈2.3垃圾输出
H10.9272 GiB96 GiB≈2.60.52 ± 0.07
H1.50.9383 GiB97 GiB≈3.40.60 ± 0.070.6672
H1.70.935OOM(×3)
H0.94OOM

窗口窄得惊人:从垃圾输出(0.80)到加载 OOM(0.935)只有 0.135 的间隔;0.92 到 0.93 的一个百分点,代价是 11 GiB/卡、回报是 +8pp GSM8K。H1.5(τlazy=0.93\tau_{\text{lazy}}=0.93、约 3.4 有效位)是全量评测的工作点:66.72% GSM8K strict-match(1319 题单 seed),注意这个数字远低于 122B 的 94.49%——3.4 位对 397B 是”可用但明显折损”的档位,论文对此没有粉饰。

H1.5 的逐类位宽分配呈现清晰的结构:100% 的自注意力、线性注意力和共享专家投影升入 xfp4;路由专家多数保持 xfp2;路由路径上的 xfp4 升级集中在第 1–3 层(GatedDeltaNet 初始过渡区)和第 13 层,且跨运行稳定。这与消融数据一致:路由专家接近高斯的主体分布可以承受激进量化,注意力不能。

吞吐:短输出 73.4 tok/s、中等 98.2 tok/s、长输出 100.9 tok/s(1500 token 输出)。一个 397B 模型以 100 tok/s 的单流速度在两张工作站卡上跑——这在 NVFP4 时代的工作站上是不可想象的。

与专家剪枝路线对比:RIY#

同一硬件约束还有另一条路线:不压位宽,砍专家。论文用 AutoRound INT4 + REAP 风格校准感知专家移除(REAP-it-Yourself 工作流)做了同模型同硬件对比:

格式专家剪枝稳态显存/卡GSM8K strict样本数tok/s(长输出)
INT4 + RIY−36%≈75 GiB33%50≈50
INT4 + RIY−24%≈80 GiB44%50≈50
XFP H10%72 GiB52%5065
XFP H1.50%83 GiB66.72%1319100.9

两条路线在不同轴上到达同一个包络:XFP+H 用每权重的位换内存、保留全部专家;RIY 用专家数量换内存、保留 4 位精度。论文的立场很克制——“都是对同一约束的合法回应”。但数据方向明确:在这个模型上,H-Process 以更低的内存(H1 的 72 GiB/卡 vs RIY-24% 的 80 GiB/卡)拿到更高的精度(52% vs 44%),H1.5 的长输出吞吐更是 RIY 的两倍(100.9 vs ≈50 tok/s),且不依赖离线校准集对专家效用的判断——REAP 的软肋恰在这里:离线删除的专家可能对某个校准集没覆盖的工作负载至关重要,而 XFP 一个专家都不丢。

硬件投影给这份工作划定了边界:H1 稳态在 2×962\times96 GB 上从容(每卡 6 GiB 余量),H1.5 勉强(余量趋零);H100 80 GB TP=2 两个剖面都放不下,H200 141 GB TP=2 有约 51 GiB 余量可从容服务 H1.5;单卡 DGX Spark(120 GB 统一内存)在当前位宽下不成立。这份工作的定位因此非常精确:工作站级、单用户、超大规模 MoE 的本地推理

讨论:作为旋钮的阈值,以及它的代价#

阈值间隔是设计出来的唯一旋钮。 在 397B 的扫描里,τlazy\tau_{\text{lazy}} 每上调 0.01,大约对应每层路由专家多一个 xfp2→xfp4 升级,稳态成本约 11 GiB/卡,回报约 +8pp GSM8K。这个可预测的换算关系让”内存 vs 质量”的权衡从试错变成查表。而在 122B 上,同一个双阈值机制自然退化回单阈值行为——N=3N=3 的中间选项吸收了弹性,宽松地板几乎不绑定。同一机制在不同模型上以不同方式生效,这个性质本身就是”质量作为约束”设计原则的验证。

位宽与离群值的收支平衡。 每个离群值 18 字节,值不值?线性路径上的均衡点:XFP3 加 y 比例离群值的开销是 0.375+18y0.375 + 18y 字节/权重,XFP4 在 2% 上限下是 0.5+0.02×18=0.860.5 + 0.02 \times 18 = 0.86 字节/权重:

0.375+18y=0.5+0.02×18y=2.7%0.375 + 18y = 0.5 + 0.02 \times 18 \quad \Longrightarrow \quad y = 2.7\%

默认 2% 上限恰好坐在均衡点附近——激进的离群值提取没有隐藏收益。这个算式也提醒读者:离群值存储是实打实的流量成本,2% 上限顶满的层,实际字节开销(0.375+0.36=0.7350.375 + 0.36 = 0.735 字节/权重,折合约 5.9 位)反而超过 INT4——所以”有效位 ≈3.97”这个数字只统计索引位宽,不含离群值残差,阅读实验数据时要记住这一点。

共享内存的代际分割。 K8192K \le 8192 的限制纯粹来自硬件分层:同一 kernel 设计在数据中心硅片上能做到 K32768K \le 32768,SM120/SM121 的每 CTA SMEM 上限(99 KB)不到数据中心 SM(228 KB)的一半,直接导致单用户解码关键维度的 4 倍回退。V2a 把上限抬回 K32768K \le 32768,代价是每通道码本精度;split-KK 是 V2 用户的另一条路,代价是 10–20% 延迟。

Lloyd 非确定性。 量化加载不产生逐位一致的权重。Qwen-35B 上相同的位宽决策在 bench.py 探针(50 题)上带来 ±8pp 的数学方差(每题 2pp,即 ±4 题的量级);权重缓存消除首次加载后的运行间方差。在 paper 级评测(1319 题 × 3 seeds)下,跨 seed 标准差收窄到 0.57%——评测样本量越大,初始化抖动被平均掉的越多。这提醒所有 XFP 用户:小样本探针上的差异不可过度解读

局限清单(论文原样列出): 单用户参考实现,并发 ≥8 时单请求吞吐崩溃(kernel 启动开销未摊销);V2 的 K8192K \le 8192 上限;无 Hessian 反馈(Hessian 加权 Lloyd 在计划中);精度评测只有 GSM8K;硬件验证限于两款工作站平台;V2a 的 N=3N=3 缺失(g=80g=80 打包路径被推迟)。最后一条与论文的核心理念形成呼应:量化格式的设计被内核打包约束反向塑造——这正是”算法-系统协同设计”的典型形态。

小结#

XFP 的价值可以浓缩成一句话:它把量化从一个”选位宽的赌博”变成一场”设质量地板的优化”。操作者用两个余弦阈值表达”我能接受多少损失”,算法用逐通道 Lloyd 码本、稀疏离群值分离和逐层自动选位去满足它,融合 kernel 把省下的字节数直接兑现成吞吐。

三个数字支撑这条路线:122B 模型 3.97 有效位、138 tok/s、GSM8K 94.49%——质量无损的速度提升;397B 模型 3.4 有效位塞进 2×962\times96 GB、100.9 tok/s、66.72%——不剪一个专家的工作站级部署;以及 H-Process 扫出来的那个 0.135 宽的阈值窗口——压缩的边界从”猜”变成了”搜”。

局限同样真实:单流场景之外的并发服务没有解决,评测面窄,Lloyd 非确定性需要权重缓存来规避。但从更大的视角看,XFP 代表的方向与 NVFP4 恰成镜像:硬件格式用固定码本换专用加速,XFP 用学习码本换硬件无关——在 NVFP4 的 Tensor Core 路径缺位的工作站层级,后者暂时跑在了前面。

参考资料#

  1. XFP: Quality-Targeted Adaptive Codebook Quantization with Sparse Outlier Separation for LLM Inference(arXiv:2605.14844,2026 年 5 月)
  2. XFP 实现代码(vLLM MultiQuant 分支)
  3. GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers
  4. AWQ: Activation-aware Weight Quantization for LLM Compression and Acceleration
  5. AutoRound: Optimize Weight Rounding via Signed Gradient Descent for the Quantization of LLMs
  6. SpQR: A Sparse-Quantized Representation for Near-Lossless LLM Weight Compression
  7. SqueezeLLM: Dense-and-Sparse Quantization
  8. AQLM: Extreme Compression of Large Language Models via Additive Quantization
  9. QuIP#: Even Better LLM Quantization with Hadamard Incoherence and Lattice Codebooks
  10. QTIP: Quantization with Trellises and Incoherence Processing
  11. REAP: Reaping Experts for Activation-Aware Pruning of MoE Models
  12. Marlin kernel(IST-DASLab)
  13. Microscaling Data Formats for Deep Learning(OCP MX 规范论文)
  14. NVIDIA Blackwell 架构页(NVFP4 相关材料)
  15. vLLM / llama.cpp / TensorRT Model Optimizer / lm-evaluation-harness

文章分享

如果这篇文章对你有帮助,欢迎分享给更多人!

XFP:质量地板驱动的自适应码本量化——3.4 位在双卡工作站跑通 397B MoE
https://pinghaoyang.com.cn/aigc/posts/xfp/
作者
平昊阳
发布于
2026-08-16
许可协议
CC BY-NC-SA 4.0

评论区

Profile Image of the Author
平昊阳
乘长风,破巨浪, 展鸿图于未央!
--
总访问量
--
访客数
公告
欢迎来到我的个人博客!欢迎关注交流吖!
更多相关公告,见
社交-留言」。
音乐
封面

音乐

暂未播放

0:000:00
暂无歌词
站点统计
文章
66
分类
16
标签
93
总字数
477,284
运行时长
0
最后活动
0 天前

文章目录