H200 · HOPPER GH100 · 特性专题

FP8 · 原生低精度 Tensor Core

★ Hopper 精度主线 · E4M3(精度/前向) + E5M2(范围/反向) · 存储减半 · 吞吐翻倍

是什么

DEFINITION

Hopper 首次在 Tensor Core 原生支持 FP8 双格式：E4M3（4 指数 + 3 尾数）与 E5M2（5 指数 + 2 尾数），并提供高速 cvt 转换。一条 MMA 把 4 个 FP8 打包进 32-bit 寄存器（k=32 vs FP16 的 k=16）——每字节吞吐翻倍。

为什么需要（vs A100 FP16/BF16）

RATIONALE · VS A100 FP16/BF16

吞吐翻倍：FP8 dense 1979 TFLOPS vs FP16 dense 990 TFLOPS（见主页 §5）；显存/寄存器占用减半。
双格式分工：单一 FP8 格式无法同时满足"前向要精度 + 反向要范围"，故定义 E4M3（前向 weights×activations）+ E5M2（反向 gradients）。这是 FP8 训练的标准配方。
vs INT8：指数自带"隐式 scale"，更能容纳 softmax / 梯度等宽动态分布，无需预设 fixed-point scale。
eager mode：运行时与 FP16/BF16 互转，per-tensor scale 量化。

格式对比（选精度的核心表）

FORMAT COMPARISON

格式	指数/尾数	bias	最大有限值	最小规格化	Inf / NaN
FP8 E4M3FN ★	4 / 3	7	448	2⁻⁶ (0.0156)	无 Inf；NaN 唯一 `S.1111.111`
FP8 E5M2	5 / 2	15	57344	2⁻¹⁴ (6.1e-5)	有 Inf `S.11111.00`；NaN `S.11111.{01,10,11}`
FP16	5 / 10	15	65504	2⁻¹⁴ (6.1e-5)	IEEE Inf/NaN
BF16	8 / 7	127	≈3.39e38	2⁻¹²⁶ (1.18e-38)	IEEE Inf/NaN

CAUTION · E4M3 没有 Inf

溢出不会变成 ±Inf，cvt.satfinite 直接饱和到 ±448。调试精度时数值贴着 448 跑，多半是被饱和——不能用"先溢出到 Inf 再判 NaN"的 IEEE 思路。

关键数字

KEY NUMBERS · SHAPES

项	值	注
FP8 wgmma 形状	m64nNk32	k=32（每寄存器打包 4 元素）
N 取值	{8,16,32,64,96,128,192,256}	同其它 dtype
累加器	f32 或 f16	默认 f32；f16 省寄存器但有溢出风险
操作数组合	4 种	e4m3×e4m3 / e4m3×e5m2 / e5m2×e4m3 / e5m2×e5m2
前向 / 反向分工	E4M3 / E5M2	前向权重激活要精度，反向梯度要范围
转换指令	cvt.rn.satfinite.e4m3x2.f32 等	成对打包；satfinite = 饱和钳位

可编译示例

COMPILABLE EXAMPLE · VERBATIM

类型转换逐字自 CUDA Math API；wgmma 助记符逐字自 CUTLASS。

#include <cuda_fp8.h>       // __nv_fp8_e4m3 / __nv_fp8_e5m2
#include <cuda_bf16.h>
#include <cuda_fp16.h>

// (a) 设备端 FP8 类型转换（构造/转换运算符，饱和语义 __NV_SATFINITE）
__device__ __nv_fp8_e4m3  f16_to_e4m3(__half h)         { return __nv_fp8_e4m3(h); }      // 前向权重
__device__ __half         e4m3_to_f16(__nv_fp8_e4m3 e)  { return static_cast<__half>(e); }
__device__ __nv_fp8_e5m2  bf16_to_e5m2(__nv_bfloat16 b){ return __nv_fp8_e5m2(b); }      // 反向梯度(宽范围)
__device__ __nv_bfloat16  e5m2_to_bf16(__nv_fp8_e5m2 e){ return static_cast<__nv_bfloat16>(e); }

// (b) FP8 wgmma 助记符（verbatim，CUTLASS mma_sm90_gmma.hpp，SS 形式）
//   wgmma.mma_async.sync.aligned.m64n8k32.f32.e4m3.e4m3  {%0..%3}, %4, %5, p, %7, %8;
//   组合: .f32.e4m3.e5m2 / .f32.e5m2.e4m3 / .f32.e5m2.e5m2 （及 .f16 累加变体）
//   N ∈ {8,16,32,64,96,128,192,256}；需 wgmma.fence / commit_group / wait_group 配套

NOTE · Transformer Engine 是什么

Transformer Engine（TE）不是 ISA 特性，而是一个库——它在每层分析张量统计、动态算 per-tensor scale、自动在 FP8 与 FP16 间切换、管理 delayed scaling 的 amax buffer。底层调用的就是上面的 FP8 wgmma。手写 FP8 kernel 时这些 scaling buffer 要自己管。

注意事项 / 陷阱

PITFALLS · CHECKLIST

CAUTION · 陷阱清单

E4M3 没有 Inf（见 §3 警告）；cvt.satfinite 溢出钳到 ±448（非 wrap / 非 Inf）。
scale factor 决定精度：per-tensor scale 常取 amax 反推；delayed scaling 用历史 amax，选错会发散/掉精度（TE 帮你管这些 buffer）。
累加器默认 FP32；FP16 累加变体省寄存器但有累加溢出风险。
eager 转换有开销：每个 FP8 GEMM 前后都要 FP8↔FP16/BF16 + 乘 scale，会吃掉一部分 FP8 收益——所以 A/B 应从 smem 喂（SS），并把转换重叠在 cp.async / wgmma 之后隐藏。

数据源

SOURCES

ONNX Float8 Technical Spec（E4M3/E5M2 位布局、bias、max、NaN/Inf 编码权威表）
CUDA Math API §15.18 __nv_fp8_e4m3（构造/转换运算符）
CUTLASS mma_sm90_gmma.hpp（FP8 wgmma 助记符）
NVIDIA Blog — Floating-Point 8（E4M3 前向 / E5M2 反向分工）