DeepSeek V4 源码剖析

第12章 FP4 e2m1 + FP8 e4m3：混合精度的几何

作者 杨艺韬 · 7,720 字 · 发布于 · 更新于

精度不是一个旋钮，是一排旋钮，每一个都要单独拧到位。

V4 的精度配方有 5 个维度：weight 精度、activation 精度、scale 精度、block 大小、计算累加精度。每个维度都要单独取舍。

12.1 引子：从 BF16 到 FP4 的精度阶梯

把 V4 涉及的浮点格式列在一张表上：

格式 总位数 符号 指数 尾数 动态范围 V4 中的位置
BF16 16 1 8 7 ~1e±38 激活默认 / KV cache（参考实现全用 BF16，部署时只有 RoPE 维用）
FP8 e4m3 8 1 4 3 2^-9 ~ 448 linear 权重 / GEMM 输入 / 部署时 KV 的非 RoPE 维
FP8 e5m2 8 1 5 2 2^-16 ~ 57344 (V3 训练时全部张量用 e4m3；V4 的推理代码与技术报告也都不涉及 e5m2)
FP4 e2m1 4 1 2 1 [-6, 6] 离散 routed expert 权重 / 部署时 Indexer 的 QK 路径
ue8m0 8 0 8 0 2^-127 ~ 2^127 scale 张量
FP32 32 1 8 23 ~1e±38 累加器 / norm 计算

V4 的精度策略是"精度阶梯化"——不同算子用不同精度，让"精度成本"最优分配。具体地：

  • 存储用极低精度（FP4 expert / FP8 linear）—— 1.6T 参数的 checkpoint 约 865 GB（64 个分片）
  • 计算用 FP32 累加—— 保证乘加结果的精度
  • 激活值默认 BF16—— 在每个算子之间保持稳定
  • scale 用 ue8m0—— 一种"纯指数"格式，记录每块的缩放因子
flowchart LR
  subgraph Storage["存储 (省内存)"]
    FP4w["expert weight: FP4 e2m1"]
    FP8w["linear weight: FP8 e4m3"]
    UE8M0["scale: ue8m0"]
    BF16Act["activation: BF16"]
  end
  subgraph Compute["计算 (保精度)"]
    Quant["act_quant: BF16 → FP8/FP4"]
    GEMM["FP8/FP4 GEMM<br/>累加器: FP32"]
    Dequant["输出: BF16"]
  end
  subgraph Numeric["数值稳定"]
    Norm["RMSNorm: FP32"]
    Soft["softmax: FP32"]
    Sink["attn_sink: FP32"]
  end
  Storage --> Compute --> Numeric

12.2 FP4 e2m1 的位级表示

FP4 e2m1 总共 4 位：

[ 1 sign | 2 exponent | 1 mantissa ]

可以表示的所有值：

二进制 数值 二进制 数值
0000 +0 1000 -0
0001 +0.5 1001 -0.5
0010 +1.0 1010 -1.0
0011 +1.5 1011 -1.5
0100 +2.0 1100 -2.0
0101 +3.0 1101 -3.0
0110 +4.0 1110 -4.0
0111 +6.0 1111 -6.0

只有 16 个编码（+0 与 -0 数值相同，实际 15 个不同值）。最大正值 6.0，最小非零正值 0.5——动态范围 [0.5, 6.0]。

FP4 的工程意义：

  • 单参数仅占 0.5 字节——若全部按 FP4 计，1.6T 参数是 800 GB（V4 用 torch.float4_e2m1fn_x2 打包，2 个 fp4 塞进 1 个字节；真实 checkpoint 是 FP4 expert + FP8 其余的混合，见第 15 章 §15.1）
  • 单参数动态范围有限（0.5 ~ 6）——必须配合 per-block scale 才能表达大数值
  • 相邻值的间距不是恒定的：0.5 → 1.0 → 1.5 之间是 0.5，1.5 → 2 → 3 是 0.5 和 1，3 → 4 → 6 是 1 和 2——每跨一个指数档间距翻倍，这正是浮点该有的样子。1 位尾数意味着舍入的相对误差上界是 2^-2 = 25%（对比 FP8 e4m3 的 3 位尾数：2^-4 = 6.25%，正好差 4 倍）。指数字段全 0 的那一档是次规格化，只有 0 和 ±0.5 两个值

V4 用 FP4 存 routed expert 权重——技术报告第 3.4 节给的理由是 MoE 专家权重是 GPU 显存占用的主要来源；精度损失则靠后训练阶段引入的 FP4 QAT 让模型去适应（不是靠"稀疏激活"摊薄，§12.9 会专门辨析）。

12.3 FP8 e4m3 的位级表示

FP8 e4m3 总共 8 位：

[ 1 sign | 4 exponent | 3 mantissa ]

指数偏置为 7，规格化数的最小正值是 2^-6 ≈ 0.0156（次规格化可以下探到 2^-9 ≈ 0.00195），最大值是 448——即 S.1111.110 = 2^8 × 1.75。注意 e4m3 的 fn 变体（PyTorch 的 torch.float8_e4m3fn，也是 V4 用的那个）没有 Inf，把 S.1111.111 留给了 NaN，所以最大值是 448 而不是 480。

V4 把 FP8 e4m3 用于 linear 权重和 GEMM 输入。具体配置（来自 config.json）：

"quantization_config": {
  "fmt": "e4m3",
  "scale_fmt": "ue8m0",
  "weight_block_size": [128, 128],
  "activation_scheme": "dynamic"
}

weight_block_size = [128, 128]：每 128×128 的权重块共用一个 ue8m0 scale。

activation_scheme = dynamic：激活值的 scale 在每次 forward 时动态计算（不是预定的）。

为什么 V4 选 e4m3 而不是 e5m2？

  • e4m3 的尾数多 1 位（3 vs 2），精度更高
  • e5m2 的指数多 1 位，动态范围更大但精度低
  • Linear 权重的数值分布相对集中（不像梯度那样跨多个数量级），精度比动态范围更重要
  • 配合 ue8m0 per-block scale，e4m3 的有限动态范围被 scale 补偿

常见的 FP8 训练做法是前向用 e4m3、反向梯度用 e5m2；V3 没有这样做，而是在全部张量上都用 e4m3，靠细粒度的分组 scale 弥补动态范围（V3 技术报告第 3.3.2 节）。V4 的推理代码同样只用 e4m3（torch.float8_e4m3fn）。

12.4 ue8m0：纯指数 scale

ue8m0 是一种特殊的 8 位浮点：

[ 8 exponent | 0 mantissa ]

无符号、无尾数、纯指数——只能表示 2 的幂：2^-127 到 2^127（编码 0xFF 留给 NaN，torch.float8_e8m0fnu 实测如此）。

ue8m0 的设计哲学：scale 的作用是"缩放数值的量级"，不需要精度——只需要"指数级"的覆盖范围。把 8 位全用来存指数，恰好覆盖 IEEE 754 float32 的指数范围。

V4 的工程意义（注意 FP8 权重与 FP4 权重的 scale 粒度完全不同，这是本章最容易读错的一点）：

  • FP8 权重：weight_block_size = [128, 128]，每 128×128 个元素共用 1 字节 scale——overhead = 1 / (128×128×1 byte) ≈ 0.006%
  • FP4 权重：inference/model.py:18 写死 fp4_block_size = 32，Linear.__init__ 里 FP4 分支的 scale 形状是 [out_features, in_features // 32]——沿 K 方向每 32 个 FP4 元素 1 字节 scale，overhead = 1 / (32 × 0.5 byte) = 6.25%，比 FP8 那条路高三个数量级
  • ue8m0 与 FP32 之间的转换是简单的位移，硬件实现极快（kernel.py 里的 fast_log2_ceil / fast_pow2 就是直接改 IEEE 754 的指数位）
  • 乘一个 2 的幂在数学上等价于给指数加一个整数，所以 ue8m0 scale 的"反量化"在不溢出的前提下不引入舍入误差；不过参考内核并不在位级做加法，而是把 scale 转成 FP32 乘到累加器上（见下图与 §12.7）
flowchart LR
  subgraph 存储["磁盘 / GPU 显存"]
    FP4Weight["FP4 e2m1: 4 bit"]
    UE8M0Scale["ue8m0 scale: 8 bit<br/>(FP4 权重: 每 32 个元素<br/>FP8 权重: 每 128×128 块)"]
    FP8Act["FP8 activation<br/>(act_quant 每 128 个元素一个 scale)"]
  end
  subgraph 加载["加载到 SMEM"]
    Cast["FP4 → FP8 类型转换<br/>(不乘 scale)"]
  end
  subgraph 计算["TensorCore + FP32 累加器"]
    GEMM["FP8 × FP8 GEMM<br/>(每个 K=32 子块)"]
    Scale["累加器 × act scale × weight scale"]
  end
  FP4Weight --> Cast
  Cast --> GEMM
  FP8Act --> GEMM
  GEMM --> Scale
  UE8M0Scale --> Scale

12.5 block-wise scale 的几何

weight_block_size = [128, 128] 的几何含义：把权重矩阵分成 128×128 的块，每块共享一个 scale。

先看一个走 FP8 那条路的权重，比如 attention 的 wkv：[dim=7168, head_dim=512]。Linear.__init__ 的 FP8 分支把 scale 形状算成 [ceil(out/128), ceil(in/128)]：

  • 行方向：512 / 128 = 4 块
  • 列方向：7168 / 128 = 56 块
  • 总块数：4 × 56 = 224 块，每块 1 字节 ue8m0 scale = 224 字节
  • 权重本体：512 × 7168 × 1 字节 (FP8) = 3,670,016 字节 ≈ 3.5 MiB
  • Scale overhead = 224 / 3670016 ≈ 0.006%——完全可以忽略

但 routed expert 不走这条路。expert 的 w1 / w3 是 [dim=7168 → inter_dim=3072] 的 FP4 权重，Linear.__init__ 的 FP4 分支用的是另一套几何——沿 K 方向每 32 个元素 1 个 scale，不分块：

  • 权重本体：[3072, 7168/2] 的 float4_e2m1fn_x2 = 3072 × 3584 字节 ≈ 10.5 MiB
  • scale：[3072, 7168/32] = [3072, 224] 的 e8m0 = 688,128 字节 = 672 KiB
  • Scale overhead = 1 / (32 × 0.5) = 6.25%

这 6.25% 是 FP4 路径的真实代价：算显存时"FP4 = 0.5 字节 / 参数"要修正成约 0.5625 字节 / 参数。

这种 per-block scale 的关键好处：不同块可以有不同的 scale 范围——某块全是大数值（scale=2^4），相邻块全是小数值（scale=2^-4）——两块独立缩放，不会互相干扰。

如果用 per-tensor scale（整个矩阵一个 scale），FP4 的有限范围 [0.5, 6] 必须覆盖矩阵中所有数值——只要矩阵里有一个异常大值，其余大部分较小的值就会被压扁到 0。分组 scale 解决的正是这个问题。

为什么 FP8 用 128、FP4 用 32？——直接的答案是标准：技术报告第 3.4 节写明 FP4 用的是 MXFP4，而 OCP MX 标准规定的就是每 32 个元素共享一个 E8M0 scale。从数值上看也说得通，两者可用的相对精度差了 4 倍：FP8 e4m3 有 3 位尾数（舍入的相对误差上界 2^-4 = 6.25%），一个 128×128 的块内即便数值跨度大一些也还撑得住；FP4 e2m1 只有 1 位尾数（上界 2^-2 = 25%）、只有 16 个编码（15 个不同值），块一大就必然有元素被压到网格边缘。所以 FP4 必须把组开得更小、用更多 scale 去换精度——代价就是上面那 6.25% 的 overhead。128 这个数还有个额外好处：与 TensorCore 的 tile 宽度（Hopper / Blackwell 的 WGMMA、tcgen05 常用 64 / 128）对齐，一个 GEMM tile 正好对应一对 scale。

12.6 V4 内部各模块的精度分配

把 V4 的精度分配画成表：

下面每一行都能在 inference/model.py 里找到对应的构造语句——凡是 Linear(...) 没显式传 dtype 的，走的就是 default_dtype，而 Transformer.__init__ 里 default_dtype = torch.float8_e4m3fn if args.dtype == "fp8" else torch.bfloat16，V4 的 inference/config.json 写的是 "dtype": "fp8"：

模块 / 张量 存储精度 计算精度 累加精度 源码依据（inference/model.py）
Embedding 权重 BF16 - - ParallelEmbedding 直接 torch.empty(...)，不走 Linear
Q/O/KV Linear 权重（wq_a/wq_b/wkv/wo_b） FP8 e4m3 FP8 GEMM FP32 Linear 默认 dtype；scale 形状 [out/128, in/128]
Attention wo_a FP8（checkpoint） BF16 einsum FP32 构造时 dtype=torch.bfloat16，但源码注释写明"wo_a is FP8 in checkpoint…using BF16 for simplicity"
Routed Expert 权重 FP4 e2m1 fp4_gemm（参考内核先把 FP4 升到 FP8 再做 FP8×FP8） FP32 expert_dtype = torch.float4_e2m1fn_x2；scale 沿 K 每 32 个元素一个
Shared Expert 权重 FP8 e4m3 FP8 GEMM FP32 Expert(args.dim, args.moe_inter_dim, swiglu_limit=...) 没传 dtype → 走默认 FP8，不是 BF16
Compressor wkv/wgate FP32 - - 显式 dtype=torch.float32（注释说 checkpoint 里是 bf16，运行时提到 fp32）
HC hc_attn_fn 等 FP32 FP32 - with set_dtype(torch.float32) 包住的 Parameter
Indexer wq_b FP8 e4m3 FP8 GEMM FP32 默认 dtype；FP4 出现在激活侧：q 与 Indexer 自己那份压缩 KV 都先过 Hadamard 旋转（rotate_activation），再做 fp4_act_quant(..., 32, True) 的伪量化，不是权重（部署时这条 QK 路径真正以 FP4 缓存和相乘，见技术报告第 3.4 节）
Indexer weights_proj BF16 BF16 GEMM FP32 显式 dtype=torch.bfloat16
LM Head BF16（checkpoint） FP32 GEMM FP32 ParallelHead.weight 是 fp32 Parameter，注释写明"lm_head in the checkpoint is stored in bf16…stored in fp32 for easier computation of logits"
Activation (默认) BF16 - - 进 GEMM 时由 linear() 调 act_quant(x, 128, ...)
KV Cache（参考实现） BF16 - - kv_cache buffer 是 bf16；nope 维按 act_quant(kv[..., :-rd], 64, ..., inplace=True) 做名义上的 FP8 伪量化后仍以 bf16 存回（按 pin 版源码字面语义，中间那次 Cast 的目标实为 BF16，数值上等于原样返回，见第 14 章 §14.6），rope 维不动
RMSNorm 内部 - FP32 - forward 里 x = x.float()
Softmax 内部 - FP32 - 防 exp 溢出
RoPE freqs_cis complex64 - - torch.polar 作用在 fp32 的角度上；位置到 1M 时最高频分量的角度约 1e6 rad，fp32 在这个量级的间距是 0.0625

读这张表能看到 V4 的精度策略："高频访问 + 数量大 → 用低精度，低频访问 + 关键路径 → 保高精度"。

384 个 routed expert 是数量最大的存储——压到 FP4 节省最多。lm_head 的 checkpoint 是 BF16，但运行时提到 FP32 再算 logits，源码注释给的理由是"for easier computation of logits"。HC / Compressor 的参数因为对精度敏感（Sinkhorn、压缩池化），直接以 FP32 参与计算。

两个最容易读错的点，这里再点一遍：shared expert 不是 BF16 而是 FP8（它和 routed expert 的区别只在 dtype 传不传，不在"量不量化"），参考实现里的 KV cache 不是 FP8 存储（FP8 只是 QAT 的伪量化，存回去仍是 bf16）。部署时则不同：技术报告第 2.3.4 节写明 KV 采用混合存储，RoPE 维用 BF16、其余维用 FP8，比纯 BF16 小近一半（1M 上下文按这个格式约 4.7 GB，见第 3 章 §3.10）。

12.7 反量化路径：从 FP4 到 BF16

V4 的 inference/model.py 里的模块级 linear() 函数处理 FP4 / FP8 / BF16 的统一调度：

def linear(x: torch.Tensor, weight: torch.Tensor, bias: Optional[torch.Tensor] = None) -> torch.Tensor:
    assert bias is None
    if weight.dtype == torch.float4_e2m1fn_x2:
        x, s = act_quant(x, block_size, scale_fmt, scale_dtype)
        return fp4_gemm(x, s, weight, weight.scale, scale_dtype)
    elif weight.dtype == torch.float8_e4m3fn:
        x, s = act_quant(x, block_size, scale_fmt, scale_dtype)
        return fp8_gemm(x, s, weight, weight.scale, scale_dtype)
    else:
        return F.linear(x, weight)

读这段代码的关键：

weight.dtype 决定路径：

  • FP4：先把 x 量化到 FP8（不是 FP4！），然后 fp4_gemm 处理"FP8 x FP4 → BF16"
  • FP8：把 x 量化到 FP8，然后 fp8_gemm 处理"FP8 x FP8 → BF16"
  • 其他：标准 BF16 GEMM

为什么 FP4 weight 配 FP8 activation？ 因为：

  1. FP4 权重反量化到 FP8 是无损的：e4m3 比 e2m1 多 2 位指数，技术报告第 3.4 节正是利用这一点，让 QAT 直接复用现有的 FP8 训练框架——FP8 激活 × FP4 权重与训练时的计算方式天然一致
  2. 官方那份参考内核里，FP4 GEMM 本身就是先把 FP4 "升采样"成 FP8 再做 FP8×FP8 的——inference/kernel.py:496 那行 B_shared[i, j] = T.Cast(FP8, T.Cast(FP32, B_fp4_shared[i, j]))，同函数的 docstring 也写明了策略是"cast FP4 to FP8 via float, then do FP8xFP8 GEMM"。既然乘法端最终都是 FP8，激活就没有再往下压的理由。（DeepGEMM 在 SM100 上另有 FP8×FP4 内核，见第 13 章——参考实现没有用它。技术报告第 1 节也说明，现有硬件上 FP4×FP8 的峰值算力与 FP8×FP8 相同，只是理论上未来硬件可以再高效 1/3；眼下 FP4 权重的收益主要是显存与访存。）

weight.scale 是个隐藏属性：

# 在 inference/model.py 的 Linear.__init__ 里（FP4 分支）
self.weight = nn.Parameter(torch.empty(out_features, in_features // 2, dtype=torch.float4_e2m1fn_x2))
scale_out_features = out_features
scale_in_features = in_features // fp4_block_size          # fp4_block_size = 32
self.weight.scale = self.scale = nn.Parameter(
    torch.empty(scale_out_features, scale_in_features, dtype=torch.float8_e8m0fnu))

V4 把 scale 直接挂到 weight 张量上（weight.scale = ...），让 linear 函数能拿到它。这是个 PyTorch 的非主流用法——通常 attribute 被挂到 Module，但 V4 直接挂到 Parameter，简化了调用。

12.8 与其他低精度方案的对比

把 V4 的精度方案与几个主流方案对比：

方案 Weight Activation Scale 适用模型
GPTQ (INT4) INT4 FP16 per-row / per-group Llama / OPT 推理
AWQ (INT4) INT4 FP16 per-group（常用 128） Llama / Qwen
FP8 (NVIDIA) FP8 e4m3 FP8 e4m3 per-tensor H100 / B200 通用
MXFP6 / MXFP4 FP6 / FP4 FP6 / FP4 per-32 OCP MX 标准，Blackwell 起原生支持
DeepSeek V3 FP8 e4m3 FP8 e4m3 权重 128×128 块 / 激活 1×128 tile（未限定 ue8m0，仅部分激活 scale 取 2 的整数幂） V3 训练 + 推理
DeepSeek V4 FP4 expert + FP8 linear FP8 ue8m0：FP8 权重 per-128×128 / FP4 权重 per-32 V4 训练 + 推理

V4 的精度方案在两个维度上独特：

  • 混合 FP4 + FP8：在 1.6T 规模的开源模型上把 FP4 专家权重与 FP8 其余权重一起落地
  • ue8m0 scale：FP4 权重这一路就是 OCP MX 标准里的 MXFP4（技术报告第 3.4 节），硬件友好

NVIDIA 从 Blackwell（B200）起已原生支持 MXFP4 / MXFP6——V4 的 FP4 权重与 OCP 标准对齐。不过按技术报告第 1 节，现有硬件上 FP4×FP8 的峰值算力与 FP8×FP8 相同，要等未来硬件才可能在算力上再提升（理论上约 1/3）；这是官方的展望，不是已实现的收益。

12.9 动手实验：用 PyTorch 模拟 FP4 量化

import torch

def fake_fp4_quant(x, block_size=32):
    """模拟 FP4 e2m1 量化 + 反量化，看精度损失"""
    # FP4 e2m1 的 16 个编码
    fp4_values = torch.tensor([0, 0.5, 1, 1.5, 2, 3, 4, 6, -0, -0.5, -1, -1.5, -2, -3, -4, -6])

    shape = x.shape
    x_flat = x.flatten()
    blocks = x_flat.split(block_size)
    out_blocks = []

    for b in blocks:
        # 找到 block 的最大绝对值，作为 scale 基准
        amax = b.abs().max()
        # ue8m0 scale = 2^ceil(log2(amax / 6))
        if amax > 0:
            scale_exp = torch.ceil(torch.log2(amax / 6.0))
            scale = torch.pow(2.0, scale_exp)
        else:
            scale = torch.tensor(1.0)

        # 量化：找最近的 FP4 值
        b_scaled = b / scale  # 现在 |b_scaled| <= 6
        # 找每个元素最近的 fp4 值
        diffs = (b_scaled.unsqueeze(-1) - fp4_values.unsqueeze(0)).abs()
        nearest = fp4_values[diffs.argmin(dim=-1)]
        # 反量化
        b_reconstructed = nearest * scale
        out_blocks.append(b_reconstructed)

    return torch.cat(out_blocks).view(shape)


# 测试
x = torch.randn(128, 128) * 3
x_fp4 = fake_fp4_quant(x, block_size=32)
err = (x - x_fp4).abs().mean().item()
relative_err = err / x.abs().mean().item()
print(f"Mean abs error: {err:.4f}")
print(f"Relative error: {relative_err*100:.2f}%")

跑这个实验会看到一个符合直觉的量级：FP4 只有 1 位尾数，舍入的相对误差上界是 2^-2 = 25%，实测的平均相对误差会明显低于这个上界（因为多数元素不落在最坏位置）——在 CPU 版 torch 2.11 上换几个随机种子跑，平均相对误差都在 11.7% 左右。

但不要把这个误差按"每 token 只见 6/384 个 expert"去稀释——那是个错误的直觉。每个被激活的 expert 内部的每一次乘加都在吃这份误差，稀疏激活减少的是参与计算的 expert 数，不是单次计算的相对误差。FP4 在 V4 上能成立，靠的是 QAT（第 14 章）让模型在后训练阶段适应这份误差（技术报告第 3.4 节：FP4 QAT 在后训练阶段引入），以及分组 scale 把误差限制在每 32 个元素之内——不是靠统计稀释。

12.10 V4 精度策略的"纵深防御"哲学

V4 的精度配方看起来复杂——5 种格式、不同 block 大小、不同模块用不同精度。这种复杂背后是一种"纵深防御" 的工程哲学：不让任何单个精度策略承担全部抗误差责任。

把这种纵深防御的层次拆开看：

第一层防御：FP4 / FP8 的 per-block scale

分组 scale 让每个组独立缩放（FP8 权重每 128×128 块、FP4 权重每 32 个元素）——某组的 outlier 不会拖累其他组。这是最基础的防御层。

第二层防御：QAT 让模型自适应量化误差

QAT 让模型在训练中（FP4 这一路是后训练阶段）"看到" FP4/FP8 的精度损失，主动学习"避开量化盲区"。这层防御覆盖了"数据分布特殊导致 per-block scale 不够"的情况。

第三层防御：FP32 累加 + BF16 中间

GEMM 内部累加用 FP32——避免低精度反复累加导致的舍入误差爆炸。结果输出 BF16——保留下游计算所需的精度。这层防御覆盖了"长链 GEMM 累积误差" 的问题。

第四层防御：HC / Sinkhorn / RMSNorm 的 FP32 计算

HC 的混合矩阵、Sinkhorn 的归一化、RMSNorm 都强制用 FP32——这些"算法关键点"不走低精度。注意 BF16 与 FP32 同为 8 位指数、动态范围相当，FP32 换来的是尾数精度（23 位对 7 位），不是更大的数值范围。这层防御兜住了"低精度让算法本身崩溃"的极端情况。

第五层防御：attn_sink / swiglu_limit / 数值 clamp

模型架构层面的"硬约束"——attn_sink 给 attention 留一个可以"不看任何位置"的出口，swiglu_limit 把 SwiGLU 的线性分量截在 [-10, 10]、gate 上限截在 10。要说明的是，swiglu_limit 的官方动机是训练稳定性而不是量化：技术报告第 4.2.3 节写明，loss spike 与 MoE 层的离群值相关，SwiGLU clamp（借鉴 gpt-oss）能消除离群值、稳定训练，Flash 和 Pro 整个训练过程都开着；而 FP4 QAT 只在后训练引入（第 3.4 节）。二者没有因果关系，clamp 不是 QAT 的一部分，只是客观上也收窄了进入低精度 GEMM 的数值范围。

5 层防御各司其职、互为补充——这是理解 V4 为什么敢在 expert 上用 FP4 的一个视角（这种分层是本专栏的归纳，不是官方的表述）。

12.11 V4 精度方案与硬件演进的对齐

V4 的精度选择不是凭空设计的——它与 NVIDIA 硬件的演进路线深度对齐。把对齐关系摆出来：

对齐 1：FP8 e4m3 与 H100 / B200 的 TensorCore 原生支持

H100（Hopper）与同代的 Ada 是 NVIDIA 首批原生支持 FP8 TensorCore 的 GPU。NVIDIA 的 FP8 标准是 e4m3 + e5m2 双格式——V4 选 e4m3（用于推理 forward）正好在 H100 / B200 的"快路径"上。如果 V4 选了非标准 FP8 格式，会失去硬件加速。

对齐 2：FP4 e2m1 与 B200 的原生支持

B200 所在的 Blackwell 是 NVIDIA 第一代原生支持 FP4 TensorCore 的架构（2024 年发布）。V4 在 2026 年发布，FP4 权重能在这代硬件上直接省显存和访存。但算力上要打个折扣：V4 的 expert GEMM 是 FP8 激活 × FP4 权重，技术报告第 1 节明说现有硬件上 FP4×FP8 的峰值算力与 FP8×FP8 相同，FP4 在算力上的收益要等未来硬件。

对齐 3：ue8m0 scale 与 OCP Microscaling 标准

OCP（Open Compute Project）2023 年发布了 Microscaling Formats（MX）v1.0 规范——定义了 MXFP4 / MXFP6 等"低精度数据 + 共享 scale" 的格式标准。V4 的 FP4 权重（e2m1 元素、沿 K 每 32 个共享一个 E8M0 scale）就是 MXFP4，技术报告第 3.4 节也直接写作 FP4 (MXFP4)；FP8 那一路的 128×128 分块则不属于 MX 规范。这意味着支持 MXFP4 的硬件（不止 NVIDIA）在数据格式上可以直接读 V4 的专家权重，但能否跑得快还取决于各家的内核实现。

对齐 4：block_size=128 与 H100 / B200 的 SMEM / TensorCore tile

H100 的 WGMMA 指令 N 维最大可到 256，128 宽的 tile 很常用；B200 的 tcgen05 同样支持 128/256 的 tile。V4 的 weight_block_size=128 与常用 tile 宽度对齐——一个 GEMM tile 对应一对 scale，便于在 kernel 里处理。

对齐 5：BF16 activation 与所有现代 GPU

BF16 是 A100 / H100 / B200 / 甚至 AMD MI300 的标准激活精度。V4 的激活默认 BF16 让模型在所有现代硬件上都"能跑"——只是 FP4 / FP8 路径在不同硬件上速度不同。

这 5 个对齐让 V4 在现有硬件上直接受益；未来硬件能否带来更多收益（例如 FP4×FP8 的算力提升）属于预期，要看硬件与内核的实际进展。

12.12 PyTorch 底层与 V4 精度的接口

V4 的 FP4 / FP8 / ue8m0 在 PyTorch 端依赖几个相对新的 dtype 与算子。把这些依赖梳理出来：

PyTorch dtype 依赖（这四个 dtype 的引入版本各不相同，其中 FP4 与 e8m0 是相当新的；官方推理代码的最低要求见下文的 inference/requirements.txt）：

  • torch.float8_e4m3fn：最早落地的一批 FP8 dtype 之一。V4 用于 linear weight 与 act_quant 输出。
  • torch.float8_e5m2：另一种 FP8 变体，V4 的推理路径不用。
  • torch.float4_e2m1fn_x2：V4 用于 routed expert weight。注意 _x2 后缀——它表示"两个 FP4 打包进 1 字节"，所以张量的最后一维是逻辑长度的一半（Linear 里写的是 torch.empty(out_features, in_features // 2, ...)）。
  • torch.float8_e8m0fnu：V4 的 ue8m0 scale 类型，OCP MX 系列格式的 scale 类型。

PyTorch 算子依赖：

  • torch.view_as_complex / torch.polar：用于 RoPE 的复指数计算
  • torch.distributed.all_reduce 等通信原语：与 V4 的并行类配合
  • nn.functional.silu / nn.functional.softplus：SwiGLU 与 sqrtsoftplus 的基础

版本要求：官方推理代码的依赖写在 inference/requirements.txt：torch>=2.10.0、transformers>=5.0.0、safetensors>=0.7.0、fast_hadamard_transform、tilelang==0.1.8。根目录 config.json 里的 "transformers_version": "4.57.1" 只是导出 config 时的 transformers 版本，不是运行要求。PyTorch 的硬门槛是"你的版本得有 torch.float4_e2m1fn_x2 和 torch.float8_e8m0fnu"——没有的话在构建模型时（Transformer.__init__ 设置 scale_dtype、Linear.__init__ 建 FP4 权重）就会直接报 AttributeError。另外 inference/kernel.py 依赖 TileLang、rotate_activation 依赖 fast_hadamard_transform，这两个也要装。

未来兼容（预测）：截至 torch 2.11 还没有 FP6 dtype（本机 dir(torch) 只有上面几种 float8 与 float4_e2m1fn_x2）；PyTorch 日后是否引入 float6_e3m2 / float6_e2m3 这类 MX 格式尚无定论。V4 的 FP4 权重本身就是 MXFP4，只要这些 dtype 保持现有语义，后续版本应能继续加载。

如果你想深入理解这些 dtype 在 PyTorch 内部如何实现（特别是 _x2 打包格式、ue8m0 的位级表示），可以参考《PyTorch 训练框架内核深度解析》第 3 章（c10 的 DType 枚举，含 Float4_e2m1fn_x2）与第 20 章（量化与混合精度，含 float8 dtype 与 torch._scaled_mm）。

12.13 V4 精度方案的"迁移到其他模型"指南

V4 的精度方案（FP4 expert + FP8 linear + ue8m0 + 128 block）是为它的特定结构设计的。如果你想把这套方案迁移到其他模型，几个关键点：

迁移点 1：weight 是否符合 block 量化假设

V4 的 weight 维度都是 128 的倍数（7168 / 3072 / 512 / 1536 / 1024）——天然 block 友好。act_quant 里那句 assert N % block_size == 0 是硬约束：维度不整除就直接断言失败。如果你的模型有 2880（= 32 × 90）这类能被 32 整除但不能被 128 整除的维度，要么把 block_size 降到能整除的值从而多付 scale 开销，要么 pad 到 128 的倍数。

迁移点 2：是否有对应规模的 expert

FP4 用在 expert 上划算，是因为专家权重是显存占用的大头（技术报告第 3.4 节），压它收益最大；这和 expert 数量多、稀疏激活能"摊薄"误差无关（§12.9）。专家规模小、专家权重占比低的模型，FP4 省下的显存有限，是否值得要结合 QAT 成本自己权衡。

迁移点 3：训练成本

V4 的 FP4 QAT 是在后训练阶段引入的（技术报告第 3.4 节）：优化器维护 FP32 主权重，前向时先量化到 FP4 再无损反量化到 FP8 参与计算。也就是说，要复刻这套方案，至少得有一段带 QAT 的训练；只做训练后量化（PTQ）直接压到 FP4，精度风险明显更大。

迁移点 4：硬件支持

H100 / B200 是 V4 精度方案的最优硬件。如果部署到 A100（无 FP8 / 无 FP4 原生），只能反量化后走 BF16 等软件路径，性能会明显下降（没有公开的实测数字）。

迁移点 5：DeepGEMM 依赖

V4 精度的高效实现依赖 DeepGEMM 这类定制内核库（官方 inference/kernel.py 则用 TileLang 写了一版参考内核）——cuBLAS 没有针对 128×128 分块 + ue8m0 scale 的原生支持（见第 13 章）。如果你的项目不能引入 DeepGEMM 作为依赖，迁移时需要自己写或者放弃这套精度方案。

总的迁移建议：V4 精度方案最适合"V4-like 模型 + H100/B200 + DeepGEMM 可用" 三个条件同时满足的场景。否则建议用更通用的 FP8 e4m3 + per-tensor scale 方案，损失部分性能换取通用性。

12.14 一段精度配方的"代价对账"

V4 的精度配方看似"白送性能"——FP4 expert + FP8 linear 让显存与访存都减小，似乎没有代价。但工程上没有免费午餐。把代价算清楚：

代价 1：训练复杂度

FP8 混合精度训练 + 后训练阶段的 FP4 QAT + Muon 优化器 + 多种 dtype 并存——训练栈的工程复杂度明显高于纯 BF16 训练，需要专门的 ML infra 团队维护。

代价 2：硬件锁定

V4 的精度配方在 H100 / B200 上最优——A100 这类没有 FP8 TensorCore 的 GPU 上要么慢、要么需要自己改内核。这种硬件锁定让"用旧硬件部署 V4" 困难。

代价 3：调试难度

低精度引入的数值问题不容易排查——某个 expert 输出异常可能是 FP4 反量化错、可能是 ue8m0 scale 错、可能是 act_quant 块大小错。每种错的症状相似但根因不同。

代价 4：fine-tune 限制

如前文（§12.13）所述，fine-tune 不能简单地从 BF16 切换到 FP4——必须保持训练时的精度配方。这限制了"用 BF16 fine-tune 然后量化部署" 这种常见 workflow。

代价 5：版本耦合

V4 的精度依赖 DeepGEMM / FlashMLA 的特定版本——这两个库升级时必须保证 V4 仍能跑。这种版本耦合让"用最新工具" 与"保持模型稳定" 之间有 trade-off。

理解这些代价让你做"V4 vs 其他模型" 的选择更清醒——不是"V4 的精度配方更先进就该用 V4"，而是"V4 适合愿意付这些代价的场景"。

如果你的项目对硬件 / 工具链 / 调试都要求灵活性，可能 BF16 dense 模型（即便能力略弱）反而是更好选择。

12.15 V4 精度配方在 inference 与 training 之间的"对称性"

V4 的精度配方有一个常被忽略的特性：inference 与 training 的精度路径高度对称。技术报告第 3.4 节把这一点当作目标来写：RL 的推理与 rollout 阶段直接用真 FP4 权重，"确保采样时的模型行为与线上部署完全一致"。

对称点 1：相同的 dtype 集合

  • weight：推理用真 FP4 expert + FP8 linear；训练时优化器维护 FP32 主权重，expert 前向先量化到 FP4、再反量化到 FP8 参与计算。FP4→FP8 的反量化是无损的（e4m3 比 e2m1 多 2 位指数，技术报告第 3.4 节），所以两边乘进 GEMM 的是同一组数值
  • activation：BF16 默认
  • scale：FP4 权重用 MXFP4 的 E8M0 scale
  • 累加器：FP32

对称点 2：相同的 act_quant 调用点

推理代码里 KV 与 Indexer 的伪量化都注明是为了对齐 QAT（inference/model.py:505 的 # FP8-simulate non-rope dims to match QAT，以及两处 # We performed QAT here），说明这些调用点是照着训练时的量化位置放的。训练代码没有公开，无法逐个核对 block_size 与 dtype 是否完全一致。

对称点 3：同一套 GEMM 库

DeepGEMM 的 README 明确它同时覆盖前向与反向（2025.05.14 起提供 dense 与 MoE 的 weight gradient kernel），也就是训练与推理共用同一份 kernel 代码。不过 V4 自己的训练脚本没有公开，"训练时具体调了哪几个 API"无法从公开材料核实——这里只能说库这一层是共用的。

对称点 4：相同的 RMSNorm / Sinkhorn

数值稳定相关的 op（RMSNorm、Sinkhorn）在推理代码里都用 FP32 计算（训练侧代码同样未公开）。

为什么这种对称性重要：

  • 零分布漂移：训练时见到的精度 = 推理时见到的精度，模型不会因部署而行为变化
  • 可复现的 bug：训练时遇到的数值问题在推理时一致出现，便于调试
  • fine-tune 的可预测性：fine-tune 用同样的精度配方，结果对部署可预测

对称性的代价：

如果你想在不同设备上部署不同精度（如 H100 用 FP8、A100 用 BF16），这种对称性会被破坏——A100 上的模型行为会与训练时有差异。要享受对称性带来的一致，部署端就得支持这套 FP8 / FP4 格式。

这是 V4 在"训推一致"与跨硬件灵活性之间的取舍，它把前者放在了更高的优先级。

12.16 延伸阅读

  • OCP Microscaling Formats（ocp-microscaling-formats）：MXFP4 / MXFP6 标准
  • FP8 Format（NVIDIA whitepaper）：H100 的 FP8 格式
  • GPTQ 论文（arXiv:2210.17323）：INT4 后训练量化
  • AWQ 论文（arXiv:2306.00978）：activation-aware 量化
  • 本专栏第 13 章：DeepGEMM 内部如何做 FP4/FP8 GEMM
  • 本专栏第 14 章：QAT 训练时的 act_quant 路径

12.17 本章小结

  • V4 用 5 维精度策略：weight (FP4/FP8) + activation (BF16/FP8) + scale (ue8m0) + block（FP8 权重 128×128、FP4 权重每 32 个元素）+ 累加 (FP32)
  • FP4 e2m1 仅 16 个编码（15 个不同值），配合 per-block scale 才能表达大动态范围
  • FP8 e4m3 比 e5m2 精度高、动态范围小——配合 scale 后整体范围能覆盖
  • ue8m0 是纯指数 scale，乘它等价于加指数、不引入舍入误差；FP4 权重这一路就是 MXFP4
  • block 128×128 与 H100/B200 的 TensorCore tile size 对齐
  • 精度阶梯化：高频大量数据用低精度（routed expert FP4），关键路径保高精度（HC / RMSNorm / softmax 用 FP32）

第 13 章我们进入 DeepGEMM——V4 这套精度配方背后的 CUDA kernel 实现。