B200// NVFP4
← B200 速查主页 / 特性卡片墙

B200 · BLACKWELL GB100 · 特性专题

NVFP4 · 原生 4-bit 低精度

★ Blackwell 精度主线 · e2m1 + 每 16 元素 E4M3 block scale · 9 PF dense 的前提

01

是什么

DEFINITION

Blackwell Tensor Core 原生支持的 4-bit 浮点格式族：基础格式 e2m1（1 符号 + 2 指数 + 1 尾数，每元素 4 bit），配合 block scaling——每 16 个元素共享一个 E4M3 (FP8) 缩放因子。这套「e2m1 + E4M3/16」组合即 NVIDIA 主推的 NVFP4；同级还有面向 MX 生态的 MXFP4（scale 用 UE8M0、每 32 元素一组）和 FP6（e3m2 / e2m3）。

为什么必须有 block scale：e2m1 裸格式只有 ±6 的动态范围、9 个可表示值，直接装权重精度会崩；per-16 缩放把每个小块自动归一到 e2m1 的甜点区，是「4-bit 也能做前向」的数值基础。

02

为什么（vs FP8）

RATIONALE · VS FP8 (HOPPER)
维度Hopper FP8 (E4M3/E5M2)Blackwell NVFP4
元素位宽8 bit4 bit (e2m1) + 0.25 bit/元素 scale 开销
动态范围裸格式自带（E4M3 ±448）靠 per-16 E4M3 scale 动态扩展
TC 吞吐 (dense)B200 4.5 PF（H200 1.98 PF）9 PF，2× FP8[1]
显存占用1×（vs FP16 0.5×）0.5× FP8——容量与带宽双省
主要用途训练 / 推理皆可推理权重+激活、部分训练场景（W4A4）
scale 路径per-tensor / per-channel（软件）block scale 硬件原生（.kind::mxf8f6f4）

两个红利叠加：算力 2×（9 vs 4.5 PF dense）与带宽 2×（4-bit 搬运量减半）。对 decode 型 memory-bound 推理，后者往往更值钱——权重减半直接把 HBM 流量砍半。

03

关键数字

KEY NUMBERS · CONSTRAINTS
项值 / 约束注
FP4 dense / sparse9,000 / 18,000 TFLOPS[1]1000W 单卡口径；GB200 10/20 PF
e2m1 可表示值9 个（±{0,.5,1,1.5,2,3,4,6}）无 scale 不可直接用
NVFP4 scaleE4M3 per 16 元素开销 8b/16elem = 0.25 bit/elem
MXFP4 scaleUE8M0 per 32 元素（power-of-2）OCP MX 标准，精度略低于 NVFP4
FP6 吞吐与 FP8 同率 4.5 PF dense[1]e3m2 / e2m3 两种，W6A6 折中
指令路径tcgen05 .kind::mxf8f6f4 + block scalescale 因子经 SMEM 旁路进 TC
有效位宽4.25 bit/元素（含 scale）vs FP8 8 bit / MXFP4 4.2 bit
04

可编译示例

COMPILABLE EXAMPLE · VERBATIM

PTX 侧走 .kind::mxf8f6f4 + scale；C++ 侧 CUDA 12.8 提供 __nv_fp4_e2m1 类型；生产级量化以 CUTLASS / TensorRT Model Optimizer 为准。

① Host 侧量化（FP32 → NVFP4 e2m1 + E4M3 scale/16）

// 每 16 个 FP32 一组：组内 max → E4M3 scale → 元素除以 scale 落 e2m1
for (size_t g = 0; g < n; g += 16) {
  float mx = fabsf(*std::max_element(w+g, w+g+16,
      [](float a, float b){ return fabsf(a) < fabsf(b); }));
  sf[g/16] = to_e4m3(mx / 6.0f);   // 让组内最大值 ≈ e2m1 满量程 6
  for (int i = 0; i < 16; ++i)
    q[g+i] = to_e2m1(w[g+i] / from_e4m3(sf[g/16])); // 2 元素/字节 pack
}

② Device 侧 tcgen05 block-scale MMA（.kind::mxf8f6f4）

// idesc 编码 e2m1×e2m1 / E4M3 scale；scale 矩阵经 SMEM 旁路送 TC
if (threadIdx.x == 0) {
  asm volatile("tcgen05.mma.cta_group::1.kind::mxf8f6f4"
                " [%0], %1, %2, %3, %4;"
    :: "r"(d_tmem), "r"(a_desc), "r"(b_desc),
       "r"(idesc_nvfp4), "r"(scale_c)); // d = (a·sfa)·(b·sfb)ᵀ + d
  asm volatile("tcgen05.commit.cta_group::1.mbarrier::arrive::one"
                ".shared::cluster.b64 [%0];" :: "r"(mbar));
}
05

注意事项 / 陷阱

PITFALLS · CHECKLIST
CAUTION · 陷阱清单
  • 9 PF 的口径前提是 NVFP4（block scale）：裸 e2m1 / per-tensor scale 达不到官方精度口径，先确认量化方案再算 roofline。
  • scale 布局有硬约束：scale 因子按 K 维 per-16 排布且需与 TC 读取布局对齐（CUTLASS 有现成 atom），自研核不同步对齐 = 静默精度崩。
  • FP4 适合前向推理与大 batch；训练中 activation 梯度、master weight 仍需高精度副本（W4A4 目前是少数场景，主流是 W4A8/W4A16 混合）。
  • NVFP4 与 MXFP4 不可混算：scale 格式（E4M3 vs UE8M0）和分组（16 vs 32）都不同，idesc 选错不报错但结果错。
  • 量化误差要按端到端评测验收（PPL / 下游任务），不要只看 MSE——4-bit 的失败模式是长尾 token 概率漂移。
  • 有效带宽红利只有权重流能吃满；KV cache 用 FP8/FP4 需 attention kernel 配合（FA4 / TensorRT-LLM 路径）。
06

数据源

SOURCES