第14章 QAT 与 act_quant：量化感知训练全链路

按上线的样子训练。推理是 FP8，训练循环里就得有 FP8。

V4 的 QAT 不是训完之后离线量化的"补救"，而是在后训练阶段就让模型带着 FP4 量化误差训练。

---

14.1 引子：PTQ vs QAT 的根本差别

把模型从 BF16 量化到 FP8 / FP4 通常有两种路径：

PTQ (Post-Training Quantization, 后训练量化)：

• 模型用 FP32 / BF16 训完
• 训完之后离线"量化校准"——根据少量样本估算每层 weight 的最佳 scale
• 推理时按校准的 scale 量化运行
• 优点：训练流程不变
• 缺点：量化引入的精度损失没有"训练补偿"——某些极敏感层会塌陷

QAT (Quantization-Aware Training, 量化感知训练)：

• 训练时在 forward 路径上插入"假量化"——把激活/权重量化再反量化，模拟推理时的精度损失
• 反向传播时让模型"感知"量化误差，自适应地学到对量化鲁棒的参数
• 优点：模型主动学会"对抗量化误差"，量化后精度损失极小
• 缺点：训练成本上升（多了 quant/dequant 操作）

V4 选 QAT 的理由：

• FP4 expert 是极端低精度，技术报告第 3.4 节的说法是引入 QAT"让模型适应量化带来的精度退化"；直接 PTQ 压到 FP4 精度风险明显更大
• DeepSeek 在 V3 上已经把 FP8 混合精度训练工业化；FP4 权重反量化到 FP8 是无损的，所以整条 QAT 管线能不加修改地复用现有 FP8 训练框架（技术报告第 3.4 节）
• 时机上，FP4 QAT 只在后训练阶段引入（技术报告第 3.4 节），预训练走的是 FP8 混合精度；作用对象是 MoE expert 权重和 Indexer 的 QK 路径

flowchart LR
  subgraph PTQ["PTQ 路径"]
    PT1[BF16 训练] --> PT2[训练完成]
    PT2 --> PT3[离线校准 scale]
    PT3 --> PT4[FP8/FP4 推理]
  end
  subgraph QAT["QAT 路径 (V4)"]
    QT1[BF16 / FP8 训练循环]
    QT1 --> QT2[forward: act_quant 注入]
    QT2 --> QT3[backward: STE 透传梯度]
    QT3 --> QT1
    QT1 --> QT4[训练完成]
    QT4 --> QT5[直接 FP8/FP4 推理]
  end
  PT4 -.精度可能塌陷.-> Bad[质量下降]
  QT5 --> Good[模型已适应量化误差]

---

14.2 act_quant 的 V4 源码

V4 在 inference/model.py 里没有 act_quant 的具体实现——它从同目录的 kernel 模块导入：

from kernel import act_quant, fp4_act_quant, fp8_gemm, fp4_gemm, sparse_attn, hc_split_sinkhorn

这里要纠正一个常见误解：act_quant 不在 DeepGEMM 里。它就在 V4 自己的 inference/kernel.py 中，用 TileLang 写成（@tilelang.jit 装饰的 act_quant_kernel）。DeepGEMM 那边反而明确说了不管这件事——README 原话是"输入转置或 FP8 casting 这类操作必须由用户单独处理"。

两个函数的真实签名（inference/kernel.py）：

def act_quant(x, block_size=128, scale_fmt=None,
              scale_dtype=torch.float32, inplace=False) -> torch.Tensor:
    """按 block 做 FP8 量化。inplace=True 时做融合的 quant+dequant，结果写回 BF16。
       scale_fmt 非空时，scale 被 round 到 2 的整数幂（MXFP 口径）。"""

def fp4_act_quant(x, block_size=32, inplace=False) -> torch.Tensor:
    """按 block 做 FP4 量化。inplace=True 时做融合的 quant+dequant。"""

三个必须看清的细节：

细节 1：分组是一维的，不是 128×128。scale 张量的形状是 [*x.shape[:-1], N // block_size]——沿最后一维每 block_size 个元素一个 scale，同一行的其他位置各算各的。那个 128×128 是权重的 scale 粒度（weight_block_size），两回事。

细节 2：inplace 决定它是真量化还是伪量化。inplace=False 时返回 (y, s)，y 是真正的 torch.float8_e4m3fn（FP4 版则是 float4_e2m1fn_x2，最后一维减半）——这条路给 GEMM 用。inplace=True 时 kernel 内部量化完立刻乘回 scale，输出 dtype 与输入相同（BF16），函数把结果 copy_ 回 x 并返回 x——这条路就是 QAT 的伪量化。

细节 3：scale 是不是 2 的幂由 scale_fmt 决定。act_quant 里的分支是 round_scale = scale_fmt is not None：为真时走 fast_round_scale（fast_pow2(fast_log2_ceil(amax * fp8_max_inv))，纯位运算求 2 的幂），为假时直接用 amax * fp8_max_inv 这个连续值。V4 的 config 里 scale_fmt = "ue8m0"，所以走的是 2 的幂那条。fp4_act_quant 则无条件用 fast_round_scale——FP4 的 scale 只能是 2 的幂。

还有一处容易忽略的下界保护：act_quant 里 amax_local[i] = T.max(amax_local[i], 1e-4)，fp4_act_quant 里是 T.max(amax_local[i], 6 * (2**-126))——防止全零块把 scale 算成 0。

---

14.3 fake quantization：训练时的"假量化"

QAT 的核心是 fake quantization——训练时做 quant + dequant 一对操作，模拟推理时的精度损失，但保持梯度流动。

具体地：

def fake_quant(x, block_size, dtype):
    """量化 + 反量化，等于训练时模拟推理量化"""
    # 1. 计算 per-block scale
    blocks = x.view(-1, block_size)
    amax = blocks.abs().max(dim=1, keepdim=True).values
    scale = compute_scale(amax, dtype)  # ue8m0 scale

    # 2. 量化
    quantized = round_to_grid(blocks / scale, dtype)  # FP4 / FP8 grid

    # 3. 反量化
    dequantized = quantized * scale

    return dequantized.view(x.shape)

forward 时这等同于"把 x 替换成 dequant(quant(x))"——数值上等于推理时的精度。

backward 时有个问题：round_to_grid 不可微（rounding 的导数几乎处处为 0，在跳变点不存在）。QAT 通常用 Straight-Through Estimator (STE)——把 round 的梯度直接 pass through：

class FakeQuant(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, block_size, dtype):
        return fake_quant(x, block_size, dtype)

    @staticmethod
    def backward(ctx, grad_output):
        return grad_output, None, None  # 直接透传梯度

这种"前向量化、后向透传"的不一致让模型感知量化的数值损失：loss 是在量化后的前向上算的，梯度则把量化当成恒等映射近似回传——两者一起让训练收敛到"对量化鲁棒"的参数。技术报告第 3.4 节描述的 V4 做法正是这样：反向对前向用的同一份 FP8 权重求梯度，直接回传给 FP32 主权重，"等价于对量化操作施加 STE"。

---

14.4 fake quantization 的 forward / backward 路径图

flowchart TB
  subgraph Forward["forward 路径（量化误差注入）"]
    F1["x: BF16"] --> F2["compute amax per block"]
    F2 --> F3["scale = 2^ceil(log2(amax/fp_max))<br/>(ue8m0)"]
    F3 --> F4["x_scaled = x / scale"]
    F4 --> F5["quantized = round_to_grid(x_scaled, FP4/FP8)"]
    F5 --> F6["dequantized = quantized × scale"]
    F6 --> F7["x_fake_quanted: BF16"]
  end
  subgraph Backward["backward 路径（STE 透传）"]
    B1["grad_output (BF16)"] --> B2["STE: 假装 round 可微"]
    B2 --> B3["grad_input = grad_output<br/>(直接透传)"]
  end
  Forward -.前向数值已被 round.-> Backward
  Backward -.后向梯度按未 round 算.-> ParamUpdate["weight 更新走<br/>'对齐量化网格的方向'"]
  classDef quant fill:#7c2d12,stroke:#fb923c
  classDef ste fill:#312e81,stroke:#a78bfa
  class F2,F3,F4,F5 quant
  class B2,B3 ste

这种"前向有损 + 后向无损" 的不对称是 STE 的关键——让模型既感知量化误差又能正常学习。

---

14.5 V4 中的 act_quant 调用点

把 V4 源码里所有 act_quant / fp4_act_quant 的调用点列出来：

位置	调用	量化目标	block_size	inplace
Attention.forward	act_quant(kv[..., :-rd], 64, scale_fmt, scale_dtype, True)	KV 的非 rope 维	64	✅ 伪量化
Indexer.forward	fp4_act_quant(q, fp4_block_size, True)	Indexer 的 query	32	✅ 伪量化
Compressor.forward（rotate=True，即 Indexer 那份）	fp4_act_quant(kv, fp4_block_size, True)	Indexer 的压缩 KV	32	✅ 伪量化
Compressor.forward（rotate=False，主 attention 那份）	act_quant(kv[..., :-rd], 64, scale_fmt, scale_dtype, True)	主 attention 的压缩 KV	64	✅ 伪量化
linear() 函数内部	act_quant(x, block_size, scale_fmt, scale_dtype)	linear 的输入激活	128	❌ 真量化

这张表最重要的信息在最后一列：前四处都是 inplace=True 的伪量化，只有 linear() 里那一处是真量化。

• 前四处的输出仍是 BF16，量化只是"让数值经历一次 FP8/FP4 的取整"——这就是 QAT 在推理代码里的残留（pin 版 FP8 那两处有个实现细节，见第 14.6 节）。源码在这几处留了注释：# FP8-simulate non-rope dims to match QAT; rope dims stay bf16 for positional precision，以及 # We performed QAT here, kv could also use fp8 format, though current implementation uses bf16
• 只有 linear() 里那次是真把激活变成 torch.float8_e4m3fn 交给 GEMM

再看粒度的分层：

• KV 与压缩 KV 走 FP8 + 64：精度敏感，且 rope 维完全不量化（kv[..., :-rd] 把最后 rope_head_dim=64 维排除在外，保住位置编码的精度）
• Indexer 那一路走 FP4 + 32：它只负责选 top-k（最近 window_size=128 个位置的滑窗 KV 不经 Indexer、始终参与注意力），可以承受更大误差；技术报告第 3.4 节也把这条 QK 路径列为 FP4 QAT 的对象之一。注意 fp4_act_quant 之前先跑了 rotate_activation——Hadamard 旋转把 outlier 摊平到各维，这是 FP4 能用的前提（见 QuaRot，arXiv:2404.00456）
• linear 输入走 FP8 + 128：与 FP8 权重的 weight_block_size=[128,128] 在 K 方向对齐

每个调用点都对应一段精度损失。FP4 expert 本身的精度损失则不在这张表里：它来自权重（FP4 权重进 fp4_gemm），按技术报告第 3.4 节，是后训练阶段对 expert 权重做 QAT 让模型适应的。

---

14.6 in-place 修饰符的语义

V4 的 act_quant 调用点大多带 inplace=True（inference/kernel.py:107 的形参名就是 inplace，没有下划线）：

act_quant(kv[..., :-rd], 64, scale_fmt, scale_dtype, True)
fp4_act_quant(q, fp4_block_size, True)

inplace=True 的语义：kernel 内部做完 quant 立刻乘回 scale，输出 dtype 与输入相同，函数再把结果 copy_ 回原张量并把它返回。FP4 那版写的是 Cast(out_dtype, Cast(FP32, Cast(FP4, clamp(x/s))) * s)，中间那次 Cast(FP4, …) 是真正的 FP4 取整。

FP8 那版有个值得注意的细节：act_quant_kernel 在 inference/kernel.py:55 先执行 out_dtype = in_dtype if inplace else out_dtype，后面伪量化那句写的是 Cast(out_dtype, Cast(FP32, Cast(out_dtype, clamp(x/s))) * s)，所以 inplace=True 时中间那次 Cast 的目标是 BF16，不是 FP8。再加上 scale 按 2 的幂向上取整，|x/s| 不会超过 448，clamp 碰不到，除以、乘以 2 的幂在 BF16 下也是精确的。结果是：pin 版 KV 与主 attention 压缩 KV 那两处"FP8 伪量化"在数值上等于原样返回。本专栏在 CPU 上按这段源码的语义用 PyTorch 复算，输出与输入逐位相同；TileLang kernel 本身未在本机实跑。注释说这里要"match QAT"，要真正模拟 FP8，中间那次 Cast 得是 FP8。

这在工程上的好处：

• 节省显存——不需要为量化后的张量分配新空间
• 让"在某段代码里 x 是 BF16，下一行就是 fake quanted BF16" 的语义直白

注意公开的这份 inference/kernel.py 是纯推理代码——它的 act_quant 里最后一句就是朴素的 x.copy_(y); return x，没有任何自定义 autograd Function。inference/model.py 里顶层的 Transformer.forward 与 MTP 模块的 forward 挂着 @torch.inference_mode()，所有子模块都在它底下运行，压根不建计算图，所以 in-place 与 autograd 的冲突在这里不存在。

训练侧要让同样的伪量化可微，就必须自己包一层 torch.autograd.Function + STE（V4 的训练代码不公开）。这也是下面"教训 4"要讲的坑：照搬这段推理代码到训练脚本里，会直接撞上 PyTorch 的 in-place 报错。

---

14.7 QAT 调试的常见问题

实现自己的 QAT 训练时容易踩的坑：

坑一：把"伪量化的粒度"和"GEMM 的 scale 粒度"混为一谈

看 §14.5 那张表会发现 block_size 有 64 / 32 / 128 三种，很容易以为"每种都要有对应的 GEMM tile"。其实不然：64 和 32 那几处根本不进 fp8_gemm / fp4_gemm——它们是 inplace=True 的伪量化，产物是 BF16。

真正要对齐的只有进 GEMM 的那条路。fp4_gemm_kernel 的文档注释把两边的粒度写得很清楚：

Act: 1x128 quant on K (reduce dim), FP8 with configurable scale dtype
Weight: 1x32 quant on K (reduce dim), FP4 with E8M0 scale

激活是每 128 个元素一个 scale，权重是每 32 个——两边故意不一样，kernel 内部按 sub_K=32 的子块处理权重、按 128 处理激活。所以"act_quant 的 block 必须等于 GEMM tile"是个错误的心智模型；正确的说法是：每种张量的 scale 粒度必须与 kernel 里对它的分块假设一致，改了一边就要改 kernel。

坑二：scale 的数据类型不匹配

V4 的 ue8m0 scale 在 PyTorch 里通常用 torch.float8_e8m0fnu dtype（如果 PyTorch 版本支持）或 torch.float32 占位。2 的幂 scale 本身用 BF16 也能精确表示（2^-126 到 2^127 都在 BF16 的指数范围内）；真正的坑在连续值 scale（scale_fmt=None 时的 amax / 448）：BF16 只有 8 位有效精度，存下来的 scale 与量化时用的 FP32 scale 不一致，反量化会带上系统性偏差。另一个坑是 scale 的 dtype 与 kernel 的约定不符——act_quant 按 scale_dtype 决定 kernel 输出 E8M0 还是 FP32，GEMM 那边读错格式就全错。

坑三：STE 的反向梯度与前向数值脱节

直接透传梯度（STE）意味着"前向被 clamp 到 ±448 的元素，反向仍按未 clamp 的梯度更新"——被截断得越狠，这个不一致越大。标准做法是配合 gradient clipping，或者对落在 clamp 区外的元素把梯度置零（clipped STE）。不过在 V4 推理代码这种"按块 amax 向上取整到 2 的幂"的 scale 下，|x/s| 不会超过 448（FP4 是 6），clamp 实际碰不到；这个坑主要出现在静态校准 scale 或可学习 scale 的方案里。V4 的训练代码不公开，具体用了哪种无从核实。

坑四：训练初期的不稳定

QAT 刚切入时会突然引入量化误差，loss 曲线通常比纯高精度训练更抖。一种常见做法是先用高精度训一段再切 QAT——V4 的 FP4 QAT 走的就是这条路：预训练用 FP8 混合精度，FP4 QAT 到后训练阶段才引入（技术报告第 3.4 节）。切入时 loss 怎么变化、有没有做预热，报告没有给。

---

14.8 与 PyTorch 的 quantization 工具对比

把 V4 的 QAT 与 PyTorch 官方 quantization 工具对比：

工具	支持精度	per-block	grouped	DeepSeek V4 兼容
torch.ao.quantization	INT8	per-tensor / per-channel	❌	❌
torch.float8_e4m3fn	FP8	per-tensor	❌	部分（需要外加 scale）
torchao FP8	FP8	per-tensor / per-row	❌	部分
V4 的 inference/kernel.py（TileLang）	FP8 e4m3 + FP4 e2m1	按最后一维分组 + ue8m0	✅	✅

V4 的 QAT 工具链与 PyTorch 主流量化工具有显著差异——V4 把 act_quant / fp4_act_quant 用 TileLang 直接写在自己的 inference/kernel.py 里，在 inference/model.py 中裸调，完全不碰 torch.ao.quantization。

这种特化让量化与 V4 自己的 kernel 贴得很紧，但也意味着：如果你想把 V4 的 QAT 思路移植到其他框架（如 transformers），需要自己实现一套 act_quant。

---

14.9 动手实验：fake_quant 模拟

import torch

def fake_quant_fp8(x, block_size=128):
    """模拟 V4 的 act_quant：FP8 e4m3 + ue8m0 per-block scale"""
    shape = x.shape
    flat = x.flatten()
    blocks = flat.view(-1, block_size)
    
    # ue8m0 scale: 2^ceil(log2(amax / fp8_max))
    fp8_max = 448.0  # FP8 e4m3 的最大值
    amax = blocks.abs().max(dim=1, keepdim=True).values
    scale_exp = torch.ceil(torch.log2((amax / fp8_max).clamp(min=1e-30)))
    scale = torch.pow(2.0, scale_exp)

    # 量化：把 blocks / scale 截断到 FP8 网格
    scaled = blocks / scale
    # 用 PyTorch 自带的 float8_e4m3fn 做真实的 FP8 舍入（就近取偶）
    quantized = scaled.clamp(-fp8_max, fp8_max).to(torch.float8_e4m3fn).float()

    # 反量化
    dequantized = quantized * scale
    return dequantized.view(shape)


# 对比原始 FP32 张量与 fake_quant 后的精度损失
torch.manual_seed(0)
x = torch.randn(2048, 2048) * 5
x_fq = fake_quant_fp8(x, block_size=128)
err = (x - x_fq).abs().mean().item()
rel_err = err / x.abs().mean().item()
print(f"Mean abs error: {err:.5f}")
print(f"Relative error: {rel_err*100:.3f}%")

# 对比矩阵乘法的精度
y = torch.randn(2048, 1024) * 5
y_fq = fake_quant_fp8(y, block_size=128)

z_bf16 = x @ y
z_qat = x_fq @ y_fq

mm_err = (z_bf16 - z_qat).abs().mean().item()
mm_rel_err = mm_err / z_bf16.abs().mean().item()
print(f"\nGEMM error after FP8 fake quant:")
print(f"Mean abs error: {mm_err:.5f}")
print(f"Relative error: {mm_rel_err*100:.3f}%")

跑这段代码会看到两件事：单张量的相对误差由 FP8 e4m3 的 3 位尾数决定（舍入的相对误差上界 6.25%，平均值明显低于此）；传到 GEMM 后，两个输入各带一份误差，K 方向上 2048 次乘加的误差又部分互相抵消，所以输出的相对误差与单张量同一量级，不会随 K 线性放大。本专栏用 CPU PyTorch 2.11.0 跑出来单张量约 2.3%、GEMM 输出约 3.8%；分布不同结果就不同，以你自己跑的为准。这就是 V4 训练时模型实际感知的"量化噪声"的量级。

---

14.10 QAT 与精度退化的攻防

V4 的 QAT 不只是"训练时模拟量化"那么简单——围绕它有一场与精度退化的"攻防战"，参与者不只是 QAT 本身。把这场攻防的几个关键回合摆出来：

攻防回合一：FP4 expert 的精度地板

FP4 e2m1 只有 16 个编码（±0 重合，实际 15 个不同值：0、±0.5、±1、±1.5、±2、±3、±4、±6）——任何 expert weight 在 FP4 下的"最低精度地板"已经被格式锁死。QAT 能做的是让模型主动避免落入 FP4 的"精度盲区"——比如让 weight 不要出现需要"在 0.5 与 1.0 之间"精确表达的中间值，而是把 weight 学得"明显偏向某一侧"。

具体做法（技术报告第 3.4 节）是：后训练阶段每次 forward 都把优化器维护的 FP32 主权重先量化到 FP4，再无损反量化到 FP8 参与计算。直觉上可以这样理解（本专栏的解读，报告没有这样描述）：模型感受到"我的 weight 实际上只有这么几个离散值"，就学会把决策边界放到这些离散点之间，避开盲区。

攻防回合二：activation 的动态范围控制

激活值在不同 layer 的数值范围可能差得很远，比如有的在 [-1, 1]，有的在 [-100, 100]。FP8 e4m3 的最大值是 448，但配合 ue8m0 scale 可以覆盖很宽的范围——前提是 scale 算得正确。

让 activation 的范围保持稳定、不要突然跨数量级，这件事在 V4 里主要不是靠 QAT，而是靠架构里的数值稳定机制：SwiGLU 的 swiglu_limit=10、HC 的 Sinkhorn 归一化等。它们客观上把 activation 的范围收窄到 FP8 友好的区间，但设计动机不是量化：技术报告第 4.2.3 节写明 SwiGLU clamp 是为消除 MoE 层离群值、稳定训练而加的，Flash 与 Pro 整个训练过程都开着；而 FP4 QAT 只在后训练引入（第 3.4 节）。swiglu_limit 不属于 QAT。

攻防回合三：scale 的离散跳级

ue8m0 scale 只能是 2 的整数幂——kernel.py 里 fast_round_scale 就是 fast_pow2(fast_log2_ceil(amax * fp8_max_inv)) 这一句（FP4 调用时传入的是 fp4_max_inv）。这是个离散决策：某个 block 的 amax 稍微越过 2 的幂，整块的 scale 就跳一级，块内所有元素的量化网格随之整体缩放一倍。

源码里能看到的是两点：一是 amax 有下界保护（FP8 那条是 max(amax, 1e-4)），避免近零块把 scale 打到极端；二是分组大小的取舍——组越小，一个 outlier 只拉低更少元素的精度，但 scale 的存储与计算开销越大。激活用 128，是与权重 weight_block_size=[128, 128] 在 K 方向对齐（第 14.5 节）；FP4 用 32，是 MXFP4 标准的组大小，也是因为它只有 1 位尾数，撑不住更大的组（第 12 章）。

攻防回合四：训练后期的"放大效应"

训练后期 lr 衰减到很小，每次更新的 weight 变化可能远小于 FP4 量化网格的间距。如果只存量化后的权重，这种小变化会"被量化网格吸收"；V4 的做法是优化器维护 FP32 主权重（技术报告第 3.4 节），小更新先在主权重里累积，越过舍入边界时才体现到 FP4 权重上，不会丢。

另一个猜测是 Muon 的正交化更新也有帮助——它把更新矩阵的奇异值拉平，更新方向更"显著"（第 17 章展开）。但 V4 没有公开这方面的训练细节，这条因果是推断而非事实。

这四个回合是理解 V4 低精度链路的几个角度；它们各自贡献多大，公开材料里没有消融数据。

---

14.11 QAT 与"训练-推理一致性"的工程承诺

V4 走 QAT 还有一个更深远的工程动机：训练-推理走同一条数值路径。

什么是训练-推理一致性？意思是：模型在训练时见到的精度、推理时跑出来的精度，走的是同一条数值路径。这听起来是个理所当然的需求，但在传统 PTQ 模型里是不成立的——PTQ 模型训练时是 BF16，推理时是 FP8，两者的数值误差从训练完成那一刻起就开始"偏移"。

V4 的 QAT 工程目标：在引入 QAT 的后训练阶段，forward 看到的量化误差与推理时一致（技术报告第 3.4 节：训练时 FP4 权重无损反量化到 FP8 参与计算，推理与 rollout 才用真 FP4，二者数值等价）。这意味着：

• 训练时看到的 loss，与线上部署后跑同样输入得到的结果，差异可控
• 你做 fine-tune 时如果改了某层 weight，可以精确预测推理时的输出变化
• A/B test 不同训练配置时，差异不会被"训练-推理精度漂移"污染

这种一致性是工程化大模型训练栈的"金标准"，做到并不容易，因为它需要训练框架与推理框架共享底层 op。V4 的公开材料指向这个方向：技术报告第 3.3 节说他们实现了端到端 batch-invariant、确定性的 kernel，目标是预训练、后训练与推理之间 bitwise 对齐；第 3.4 节说 RL 的 rollout 与推理阶段直接用真 FP4 权重，保证采样行为与线上部署一致；DeepGEMM 同时提供前向与 weight-gradient kernel；推理代码里那几处 inplace=True 的伪量化和注释（"We performed QAT here"）也说明推理路径是照着训练时的量化位置放的。但训练栈本身不公开，严格意义上的 bit 级一致无法从外部验证。

如果你的项目想从 V4 借鉴 QAT 思路，最关键的不是 act_quant 函数本身——而是确保训练和推理的 op 实现完全一致。这是 V4 最难复制的工程能力之一。

---

14.12 QAT 与 LLM 评估的"精度漂移"检测

V4 的 QAT 让训练-推理精度对齐，但生产中仍需持续监控"是否真的对齐"。这部分监控属于评估工程的范畴——与《LLM 评测体系》专栏有直接关联。

精度漂移的来源：

• 软件层：vLLM / DeepGEMM 升级时算子实现可能有微小差异
• 硬件层：不同 GPU 架构走的 kernel 实现不同（如 Hopper 与 Blackwell），数值细节可能略有差异
• 数据层：用户输入分布与训练分布不同，激活值范围偏移

漂移检测方法：

• golden test：选 100-1000 个 fixed prompt + reference output，定期跑一遍对比 token-by-token 差异
• distribution test：监控生产输出的统计分布（token 频率、长度、log-prob），与训练时 reference 对比
• shadow test：把一小部分流量同时跑 V4 production 和高精度 reference 实例，对比差异

当检测到漂移时的应对：

• 小漂移：通常无害，不必处理
• 中漂移：可能是软件版本不兼容，回滚或修复
• 大漂移：必须排查根本原因——可能是部署配置错（如精度模式被误设）

大、中、小的阈值要在你自己的模型与流量上标定，没有可以照抄的数。

这套精度漂移检测是低精度模型上线后的标配。具体的 evaluation pipeline 实现细节（指标定义、自动化跑批、告警阈值）见《LLM 评测体系》专栏第 18 章"把 Evals 变成 Quality Gate"。

---

14.13 QAT 训练栈的 5 个工程教训

V4 的 QAT 把"训练时模拟推理量化"推到了 1.6T 规模。结合公开源码与技术报告，可以归纳出几条工程教训，想复刻 QAT 的项目可以参考：

教训 1：QAT 不必从训练第一步就启用

V4 的 FP4 QAT 只在后训练阶段引入（技术报告第 3.4 节），预训练走的是 FP8 混合精度框架。这条路能走通的一个前提是：FP4 反量化到 FP8 无损，QAT 阶段可以原样复用预训练的 FP8 训练框架，不必另起一套。"从头 QAT"与"后期切入"孰优孰劣，公开材料里没有对照数据。

教训 2：scale 不需要可微

ue8m0 scale 是离散的（2 的整数幂），ceil 本身不可微，但这不妨碍训练：技术报告第 3.4 节写明，反向对前向用的同一份 FP8 权重求梯度，直接回传给 FP32 主权重，等价于对整个量化操作做 STE，scale 作为量化操作的一部分被一并"跳过"。真正要避免的是反向链路意外穿过 round/ceil 的真实导数（几乎处处为 0），那样梯度才会被截断。

教训 3：组边界的 scale 跳变

相邻两组用各自的 scale。如果两组的 amax 差出好几个数量级，边界处就会出现"同样大小的数值，落在完全不同的量化网格上"的不连续。V4 源码里能看到的直接对策是 Hadamard 旋转（rotate_activation）——它把 outlier 摊到所有维度上，让各组的 amax 彼此接近，这也是 FP4 那条路的前置步骤。

（至于"相邻块 scale 差异通常在什么范围"——这需要拿到真实权重统计才能说，本专栏没有做过这个测量，不给数字。）

教训 4：照搬推理版的 in-place 伪量化会在训练里炸

inference/kernel.py 里的 act_quant(..., inplace=True) 最后是朴素的 x.copy_(y)，没有任何 autograd 包装——因为整个推理路径跑在 @torch.inference_mode() 下，不建图。搬到训练脚本里，PyTorch 的 autograd 会因为"需要梯度的张量被原地修改"直接报错。要在训练里用同样的伪量化，必须自己包一层 torch.autograd.Function（forward 调 kernel、backward 透传），或者干脆用非 in-place 版本。

教训 5：精度漂移监控比"训练 loss 监控"更重要

QAT 的 loss 可能看起来正常下降，但前向输出与 BF16 reference 的偏差在悄悄增大。合理的做法是每隔 N step 跑同一组固定输入，对比 QAT 输出与 BF16 reference 输出的分布差异——这个指标反映"QAT 是否真的在学量化鲁棒性"，比 loss 更直接。V4 训练栈不公开，这一条是通用工程建议而非对 V4 的描述。

这 5 条中，教训 1、2 有技术报告佐证，教训 3、4 来自公开源码，教训 5 是通用工程建议。

---

14.14 QAT 工程师速记

关键概念：

• QAT = 训练时模拟推理的量化误差
• act_quant：把激活动态量化到 FP8 / FP4
• STE：Straight-Through Estimator，反向梯度透传
• in-place fake quant：节省显存的量化-反量化对

V4 中 act_quant 调用点：

• KV 的非 rope 维：FP8 + block 64（伪量化，rope 维不动）
• 主 attention 的压缩 KV：FP8 + block 64（伪量化）
• Indexer query：FP4 + block 32（伪量化，前面先跑 Hadamard 旋转）
• Indexer 的 Compressor（rotate=True）：FP4 + block 32（伪量化）
• linear 输入：FP8 + block 128（唯一的真量化，产物直接进 GEMM）

调试 QAT 的 5 个症状：

1. loss 早期 NaN —— scale 计算溢出
2. 训练曲线"阶梯感" —— scale 跳级，正常
3. 推理输出与训练 reference 偏差大 —— QAT 没真正生效，回查 act_quant 是否启用
4. fine-tune 后量化误差变大 —— fine-tune 时关掉了 QAT
5. 某些层精度异常 —— 该层的 scale 粒度与 kernel 的分块假设不一致

与 PyTorch 官方 quantization 的差异：

• V4 的 act_quant 是自己用 TileLang 写在 inference/kernel.py 里的，不用 torch.ao.quantization，也不来自 DeepGEMM
• 分组 ue8m0 scale（激活沿最后一维分组），不是 per-tensor
• 伪量化走 in-place（推理代码在 inference_mode 下不建图；训练侧要自己包 autograd Function）

生产监控指标：

• 推理输出与 reference 实例的分布差异（阈值必须在你自己的模型与流量上标定，没有可以照抄的数）
• 每层 act_quant 的 amax 分布（应稳定，无突跳）
• 部署后的精度漂移（定期跑 golden test）

---

14.15 QAT 与"反向传播误差累积" 的边界

V4-Pro 是 1.6T 参数、预训练 33T tokens（技术报告口径）的规模，低精度训练的误差累积是个真实问题。把这个边界讲清楚。

误差累积来源 1：STE 的不精确

Straight-Through Estimator 把 round 的反向梯度透传——前向数值已经被 round，反向梯度按"未 round" 算。这个不一致让某些极端方向的梯度估计有偏差。

模型越大、QAT 训练越长，这类偏差累积的机会越多；具体影响多大，公开材料里没有数据。

误差累积来源 2：FP4 / FP8 GEMM 的累加误差

虽然 GEMM 内部用 FP32 累加，但沿 K 维累加几千次乘积后，FP32 自身的舍入误差也会累积。

误差累积来源 3：跨 layer 误差放大

第 1 层的小误差会逐层传下去，61 层之后是个连乘。放大还是收敛，取决于每层的等效增益：大于 1 就指数放大，小于 1 就衰减。V4 的每层都有 RMSNorm 把 hidden 重新归一，HC 的 Sinkhorn 又约束了 4 路之间的能量分配——这两个机制正是把等效增益按在 1 附近的东西。（具体放大倍数需要实测，不是能靠公式估出来的。）

V4 架构里的 4 个数值稳定机制（它们不是为 QAT 设计的，但客观上压住了误差累积）：

1. HC 的 Sinkhorn 双随机性：双随机矩阵谱范数不超过 1，4 路 hidden 的混合不会放大信号——抑制误差累积
2. per-layer 精度复位：每层 RMSNorm 把 hidden 重新归一——切断误差链
3. attn_sink：每个头可学习的 sink 项进入 softmax 分母，允许注意力权重总和小于 1（技术报告第 2.3.3 节），选中的 KV 都不相关时不必强行分配注意力
4. swiglu_limit 的 clip：防止 SwiGLU 输出爆炸（技术报告第 4.2.3 节：为消除离群值、稳定训练而加，全程开启，不属于 QAT）

这 4 个机制一起把误差累积按在可接受范围内；各自贡献多大，公开材料里没有消融。（"和 BF16 reference 差多少"这类数字需要拿到两份权重实测，公开材料里也没有。）

理解这个边界的意义：

如果你想从 V4 借鉴 QAT 设计到自己的模型，要同时考虑架构层的数值稳定——光有 act_quant 不够，还要有配套的架构与数值兜底。这 4 个机制哪些是必需的，要在你自己的模型上验证。

---

14.16 延伸阅读

• 《A Survey of Quantization Methods for Efficient Neural Network Inference》（Gholami 等，arXiv:2103.13630）：量化方法综述，含 QAT 与 PTQ 的系统对比
• 《Estimating or Propagating Gradients Through Stochastic Neurons for Conditional Computation》（Bengio 等，arXiv:1308.3432）：STE 的源头
• QuaRot（arXiv:2404.00456）：Hadamard 旋转消 outlier——V4 的 rotate_activation 走的就是这条路
• DeepSeek-V4 技术报告第 3.4 节（FP4 Quantization-Aware Training）
• DeepSeek-V3 报告中关于 FP8 训练的章节
• 本专栏第 12 章：FP4 / FP8 / ue8m0 格式细节
• 本专栏第 13 章：DeepGEMM 怎么吃「已经量化好的 FP8 激活 + scale」（它不负责量化本身）
• 本专栏第 17 章：Muon 优化器如何与 QAT 配合

---

14.17 本章小结

• V4 用 QAT 而非 PTQ——后训练阶段就让模型"看到" FP4 量化误差（技术报告第 3.4 节）
• 推理代码里 inplace=True 的 act_quant / fp4_act_quant 在 forward 路径上复刻"假量化"（量化 + 反量化）；训练侧 backward 用 STE 透传梯度
• V4 源码里 5 个 act_quant 调用点：KV、Indexer 的 Q、Compressor 双路径、linear 输入——前四处是 inplace=True 的伪量化（输出仍是 BF16），只有 linear 输入那处是真量化
• 激活的 scale 沿最后一维分组（128 / 64 / 32），权重的 scale 是另一套粒度（FP8 权重 128×128、FP4 权重沿 K 每 32 个）——两者不必相等，fp4_gemm_kernel 就是按 128（激活）配 32（权重）写的
• QAT 调试 4 个常见坑：混淆伪量化与真量化的粒度、scale 类型错、STE 与前向 clamp 脱节、初期不稳定
• V4 的 QAT 工具链与 PyTorch 官方 quantization 工具差异显著——act_quant 是 V4 自己用 TileLang 写在 inference/kernel.py 里的，不来自 DeepGEMM

第 15 章我们离开精度链路，进入 V4 的分布式工程：Tensor / Expert 并行的实现细节。