第13章 DeepGEMM：FP4 / FP8 矩阵乘法内核

“A great GEMM kernel is invisible in your model code, but it’s the difference between an idea and a product.” —— 引自一位 cuBLAS 老兵

V4 的 fp4_gemm(x, s, weight, weight.scale, scale_dtype) 一行调用，背后是 DeepGEMM 把 H100 / B200 的 TensorCore 推到 95% 利用率的几千行 CUDA。

13.1 引子：为什么 cuBLAS 不够用

NVIDIA 的 cuBLAS 是矩阵乘法的事实标准——FP16 / BF16 / FP8 都有官方实现。为什么 V4 不直接用 cuBLAS，要自己写一个 DeepGEMM？

理由有四个：

理由一：cuBLAS 的 FP8 支持不够灵活

cuBLAS 的 FP8 GEMM 假设 scale 是 per-tensor 或 per-row 的简单形式。V4 用的是 per-block 128×128 + ue8m0 scale——cuBLAS 没有原生支持。

理由二：cuBLAS 不支持 FP4

截至 2026 年，cuBLAS 还没有 FP4 GEMM 的官方接口。V4 的 routed expert 必须自己写 FP4 GEMM。

理由三：MoE 的稀疏 dispatch

cuBLAS 是为”全 batch 一起算”设计的。V4 的 MoE 是”每个 token 选 6 个 expert”——每个 expert 看到的 token 数不同（per-batch 不规则）。cuBLAS 无法高效处理这种 grouped GEMM。

理由四：V4 的特化优化空间

为 V4 量身定做的 GEMM 可以做特化优化：与 act_quant 融合、与 RMSNorm 融合、与 HC 的 Sinkhorn 融合——cuBLAS 是通用的，没有这种”上下游融合”的能力。

DeepGEMM 仓库（github.com/deepseek-ai/DeepGEMM）的诞生就是为了解决这四个问题。

13.2 DeepGEMM 仓库结构

DeepGEMM/
├── csrc/
│   ├── deep_gemm/
│   │   ├── fp8_gemm_sm90.cu       # H100 / H800 路径
│   │   ├── fp8_gemm_sm100.cu      # B200 路径
│   │   ├── fp4_gemm_sm90.cu        # H100 FP4 (模拟)
│   │   ├── fp4_gemm_sm100.cu       # B200 原生 FP4
│   │   ├── grouped_gemm_*.cu      # MoE 用的 grouped GEMM
│   │   ├── act_quant.cu           # 与 GEMM 融合的激活量化
│   │   └── ...
│   └── deep_gemm_extension.cpp    # PyTorch binding
├── deep_gemm/                      # Python 包装
└── tests/

DeepGEMM 总代码量约 600 行核心 + 几千行 wrapper / tests。仓库的设计哲学：每个硬件架构一个独立 .cu，不抽象——避免”通用化”导致的性能损失。

13.3 fp8_gemm 的接口与算法

V4 的 fp8_gemm 接口大致是：

torch::Tensor fp8_gemm(
    torch::Tensor x_fp8,         // [M, K] FP8 e4m3
    torch::Tensor x_scale,        // [M / block_size, K / block_size] ue8m0
    torch::Tensor w_fp8,          // [N, K] FP8 e4m3
    torch::Tensor w_scale,        // [N / block_size, K / block_size] ue8m0
    torch::Tensor scale_dtype     // 输出 scale 的 dtype
);  // 返回: [M, N] BF16

算法核心：

对每个输出 tile (m_tile, n_tile):
    accumulator = 0  (FP32)
    对每个 k_tile:
        # 1. 从 global → SMEM 加载 x[m_tile, k_tile] 和 w[n_tile, k_tile] (FP8)
        # 2. 加载对应的 scale (ue8m0)
        # 3. WGMMA 指令做 FP8 × FP8 → FP32 累加
        accumulator += scale_x * scale_w * (x_block @ w_block.T)
    # 4. 把 accumulator (FP32) 转 BF16 写回
    output[m_tile, n_tile] = bfloat16(accumulator)

关键工程点：

点 1：scale 在累加器上做

每个 (m, k) 块和 (n, k) 块各有一个 ue8m0 scale。WGMMA 计算 x_block @ w_block 后，用 2^(scale_x + scale_w) 缩放结果——因为 ue8m0 是纯指数，“加法 + pow2”等价于乘法。

这种 scale 模型让 SMEM 压力很小——每 128×128 的块只多 1 字节 scale，不影响 SMEM 用量。

点 2：tile 大小与 block_size 对齐

block_size = 128（来自 V4 的 weight_block_size）。WGMMA 的 tile size 通常也是 64 或 128——让 GEMM 的 tile 与量化 block 对齐，每个 tile 用一对 scale。

点 3：异步加载 + 计算重叠

H100 / B200 的 cp.async.bulk + TMA 让全局内存到 SMEM 的拷贝异步进行——计算前一个 tile 的同时拷贝下一个 tile 的数据。这是 DeepGEMM 接近硬件极限的关键。

13.3·补 fp8_gemm 的异步流水线图

把 §13.3 的算法用图表达——展示异步加载与计算重叠：

flowchart LR
  subgraph Cycle1["Cycle 1"]
    A1["cp.async.bulk<br/>load tile 0 → SMEM"]
  end
  subgraph Cycle2["Cycle 2"]
    A2["cp.async.bulk<br/>load tile 1"]
    W1["wait tile 0"]
  end
  subgraph Cycle3["Cycle 3"]
    A3["cp.async.bulk<br/>load tile 2"]
    W2["wait tile 1"]
    C1["WGMMA compute<br/>tile 0 → FP32 acc"]
  end
  subgraph Cycle4["Cycle 4"]
    A4["load tile 3"]
    W3["wait tile 2"]
    C2["compute tile 1"]
  end
  subgraph Output["...output stage"]
    Acc["FP32 累加器"]
    Cast["→ BF16 输出"]
  end

  Cycle1 --> Cycle2 --> Cycle3 --> Cycle4 --> Output
  C1 -.scale_x × scale_w.-> Acc
  C2 -.scale_x × scale_w.-> Acc
  Acc --> Cast

每个 cycle 同时做 3 件事：加载下一个 tile + 等上一个 tile + 算上上个 tile。这种”3-stage 流水线”让 SMEM 持续被使用、TensorCore 持续工作——接近硬件极限。

如果某个 stage 落后（如 SMEM 不够大、cp.async 延迟过长），整条流水线会”气泡化”——TFlops 大幅下降。这是 DeepGEMM 调优时最重要的诊断点。

13.4 fp4_gemm 的特殊处理

FP4 GEMM 在 H100 上没有原生硬件指令——必须模拟。DeepGEMM 在 SM90 上的 fp4_gemm 路径：

1. 把 FP4 weight 反量化到 FP8 (用 ue8m0 scale)
2. 走标准 FP8 GEMM
3. 输出 BF16

代价是反量化步骤——每次 GEMM 多一次 SMEM 内的位移操作。但因为 FP4 weight 在 global memory 占用减半，整体带宽节省超过反量化开销——FP4 路径的实际吞吐与 FP8 接近。

在 SM100 (B200) 上，FP4 是原生 TensorCore 指令——不需要反量化模拟。这让 V4 在 B200 上的 expert GEMM 比 H100 快约 1.6 倍（README 公开数字）。

13.5 grouped GEMM：MoE 的关键 kernel

V4 的 MoE forward 用循环逐 expert 跑：

for i in range(self.experts_start_idx, self.experts_end_idx):
    if counts[i] == 0:
        continue
    expert = self.experts[i]
    idx, top = torch.where(indices == i)
    y[idx] += expert(x[idx], weights[idx, top, None])

这个循环在小 batch 下效率很低——每个 expert 只算 1-2 个 token，GEMM 的 launch overhead 主导计算。

DeepGEMM 提供 grouped_gemm kernel：把 N 个 expert 的 GEMM 合并到一次 launch：

torch::Tensor grouped_gemm_fp8(
    torch::Tensor x_fp8,            // [total_M, K]
    torch::Tensor w_fp8,            // [N_experts, N, K]
    torch::Tensor expert_offsets    // [N_experts + 1] 每个 expert 的 token 范围
);

输入：x 是所有 expert 输入 token 的拼接（按 expert 分组），expert_offsets 标识每段属于哪个 expert。 输出：每段对应 expert 的输出。

这种 grouped GEMM 把 N 个独立 GEMM 合并成一次 kernel launch——大幅降低 launch overhead，对小 batch MoE 至关重要。

V4 在生产部署时（vLLM 等）会优先走 grouped_gemm 路径——不是 Python 循环。inference/model.py 里的 Python 循环主要给”reference 实现”用，验证 grouped_gemm 正确性。

13.5·补 grouped_gemm 在 MoE 中的调度图

把 grouped_gemm 在 V4 MoE forward 中的位置画清楚：

flowchart TB
  Input["x: [B*S, 7168] BF16"] --> Gate
  Gate --> TopK["topk indices: [B*S, 6]"]
  TopK --> Sort["按 expert id 排序<br/>token (DeepEP dispatch)"]
  Sort --> Grouped["x_grouped: [total_tokens, 7168]<br/>+ offsets: [n_experts+1]"]
  Grouped --> GG["grouped_gemm_fp8 / fp4<br/>1 次 launch 跑 N 个 expert"]
  GG --> ExpertOuts["expert outputs"]
  ExpertOuts --> Combine["DeepEP combine<br/>按 weight 聚合"]
  Combine --> Y["y: [B*S, 7168]"]
  
  classDef gemm fill:#312e81,stroke:#a78bfa,color:#ede9fe
  class GG gemm

关键点：grouped_gemm 一次 launch 处理所有 expert 的 GEMM——不需要 N 次 kernel launch。这是 MoE 在小 batch 下不被 launch overhead 拖死的根本工程优化。

13.6 与 cutlass / cuBLAS / Triton 的对比

把 DeepGEMM 与同类方案对比：

DeepGEMM 的定位是”为 DeepSeek V3/V4 量身定制的 GEMM”——不追求通用，但在自己的特化场景下追求极致性能。

如果你想在自己的项目里用 V4 类似的 FP4 + ue8m0 + per-block 128 配方，DeepGEMM 几乎是唯一开箱即用的选择——cuBLAS 不支持 FP4，cutlass / Triton 需要你自己组合实现。

13.7 编译与部署

DeepGEMM 的编译需要：

# 必须 CUDA 12.8+ 和 GCC 11+
git clone https://github.com/deepseek-ai/DeepGEMM.git
cd DeepGEMM

# 自动检测 GPU 架构（默认编 sm_90 + sm_100）
pip install -e .

# 验证安装
python -c "from deep_gemm import fp8_gemm; print('OK')"

# 跑性能基准
python benchmark/bench_fp8_gemm.py --m 4096 --n 4096 --k 4096

集成到 vLLM / SGLang 时，通常作为 vllm 的可选依赖：

pip install vllm[deepseek-v4]

这个 extras 会自动安装 DeepGEMM + FlashMLA 两个仓库的最新版本。

13.8 性能数字（README 公开）

DeepGEMM 在 H100 上的 FP8 GEMM 吞吐（来自 README 与公开 benchmark）：

• M=N=K=4096：约 1300 TFlops（接近 H100 峰值的 90%）
• M=N=K=8192：约 1400 TFlops

cuBLAS 同等配置约 900-1000 TFlops——DeepGEMM 比 cuBLAS 快 30-40%。差距主要来自 per-block scale 的高效处理 + 与 act_quant 的融合。

V4 用 DeepGEMM 跑 GEMM，每个 forward 节省的时间累积起来，让 V4 的 token 吞吐比”用 cuBLAS 实现的同等模型”快 1.3-1.5 倍。

13.9 动手实验：跑 DeepGEMM benchmark

git clone https://github.com/deepseek-ai/DeepGEMM.git
cd DeepGEMM
pip install -e .

# 跑 V4 典型 GEMM 形状的基准
python -c "
import torch
from deep_gemm import fp8_gemm

M, N, K = 4096, 7168, 7168
block_size = 128

# 随机生成 FP8 张量 + ue8m0 scale
x = torch.randn(M, K, dtype=torch.bfloat16, device='cuda').to(torch.float8_e4m3fn)
x_scale = torch.randn(M // block_size, K // block_size, dtype=torch.float32, device='cuda')
w = torch.randn(N, K, dtype=torch.bfloat16, device='cuda').to(torch.float8_e4m3fn)
w_scale = torch.randn(N // block_size, K // block_size, dtype=torch.float32, device='cuda')

# 跑 GEMM
import time
torch.cuda.synchronize()
start = time.perf_counter()
for _ in range(10):
    y = fp8_gemm(x, x_scale, w, w_scale)
torch.cuda.synchronize()
elapsed = (time.perf_counter() - start) / 10
tflops = 2 * M * N * K / elapsed / 1e12
print(f'FP8 GEMM {M}x{N}x{K}: {elapsed*1000:.2f} ms, {tflops:.1f} TFlops')
"

跑完会得到一个具体的 TFlops 数字。如果你的 H100 / B200 跑出 < 1000 TFlops，可能是 thermal throttling 或 memory clock 设置问题——参考 NVIDIA 的 GPU 性能调优文档。

13.9·补 DeepGEMM 设计哲学的几个潜规则

DeepGEMM 的代码读起来与 cutlass、Triton 这些通用 GEMM 库非常不同。差异不在算法本身，而在工程哲学。把它的几条潜规则归纳出来：

潜规则一：拒绝模板抽象

cutlass 是 C++ 模板地狱——一个 GEMM 实现可能展开成几十层模板。DeepGEMM 反其道而行：每个硬件架构一个独立的 .cu 文件，用最直接的 CUDA C 写法。读者打开 fp8_gemm_sm90.cu 看到的就是 SM90 的具体优化，不需要在模板里”展开演算”。

这种反直觉的选择背后有具体的工程理由：V4 的 GEMM 形状是固定的（M/N/K 都是 128 的整数倍、weight_block 固定 128×128、scale 固定 ue8m0）。在固定输入下，模板的”灵活性”反而是负担——它让编译器无法做最激进的特化优化。

潜规则二：与上下游 op 主动融合

cuBLAS 把 GEMM 当作”独立黑盒”——输入张量、输出张量、不关心上下游。DeepGEMM 主动把 GEMM 与 act_quant、与 RMSNorm、与 MoE 的 token gather 融合。这种融合在算子边界节省了多次 SMEM/global memory 往返。

例如 linear 调用前必须 act_quant——DeepGEMM 的 fp8_gemm 直接接收”已量化的 FP8 + scale” 作为输入，act_quant 在 GEMM 启动前的同一 kernel 里完成。这种融合让 V4 的 forward 比”分别调用 act_quant 和 GEMM” 快约 15-20%。

潜规则三：硬编码的 tile size

cutlass / Triton 让你选 tile size。DeepGEMM 把 tile size 硬编码成 SM90 / SM100 的最优值（典型 128×128×64），不让用户改。这削弱了灵活性，但保证了”DeepGEMM 跑出的数字就是这个硬件的最优”。

潜规则四：不支持反向传播（in inference 仓库）

公开的 DeepGEMM 是”推理优先”——只支持 forward 路径，没有 backward。训练时的 backward GEMM 是 V4 团队内部的另一份代码，不公开。这种”训推分离” 让公开仓库保持极简。

这四条潜规则让 DeepGEMM 不能直接替代 cuBLAS，但作为 V4 这种特化场景的 GEMM 引擎，它的极致性能完全合理。如果你的项目想从 DeepGEMM 借鉴某个优化，要做好”无法直接复用、必须重写到自己代码里”的心理准备。

13.9·补·补 DeepGEMM 与 V4 训练栈的串接

DeepGEMM 不只是给推理用——V4 训练时同样依赖它。具体串接点：

前向 GEMM：训练每 step 的 forward 走 DeepGEMM 的 fp8_gemm / fp4_gemm，与推理路径一致。这保证了 QAT 训练时模型”看到的精度”与推理时完全相同——避免了”训练用一种精度、推理用另一种” 的常见坑。

反向 GEMM：训练的 backward 也需要 GEMM——但 backward 的 weight grad 计算 (dY @ X^T) 与 forward GEMM 的形状不同。V4 团队为 backward 写了独立的 kernel（不公开），但精度策略一致——都是 FP8 输入 + FP32 累加 + BF16 输出。

优化器步：Muon 优化器的 Newton-Schulz 迭代涉及 GEMM——DeepGEMM 提供了 BF16 GEMM 用于这一步（不需要量化，因为 Newton-Schulz 在 FP32 中算）。

蒸馏阶段：on-policy 蒸馏阶段，teacher / student 的 forward 都用 DeepGEMM——保证两者数值精度一致。这避免了”teacher 的精度比 student 高所以学不像”的隐性陷阱。

DeepGEMM 在 V4 全生命周期（预训练 / 后训练 / 推理）都是底层基础设施。理解 DeepGEMM 等于理解 V4 在硬件层的”血液系统”。

13.9·延展 DeepGEMM 性能调优经验

如果你部署 V4 后发现 DeepGEMM 没跑到预期吞吐，常见原因：

调优点 1：GPU 时钟与温度

H100 在 boost clock 下能跑 1300+ TFlops，但只有”温度 < 75°C 且功耗 < 700W”时才维持。如果你的 GPU 散热不足、或者 nvidia-smi 看到 power limit 触顶，吞吐会下降 15-25%。先检查 GPU 物理状态再调软件。

调优点 2：CUDA Stream 配置

DeepGEMM 默认在当前 default stream 跑。如果 PyTorch 的其他算子也在 default stream，会发生 stream serial（串行）。建议把 DeepGEMM 放在独立 stream，与其他算子并行。

调优点 3：SMEM bank conflict

block_size=128 与 SMEM bank 数（32 个 bank）的对齐影响很大。如果 weight 的 stride 不是 128 的倍数，会触发 bank conflict，吞吐降 20-30%。V4 的 config 已经保证所有维度都是 128 倍数——但你自定义的 fine-tune 模型可能违反这个对齐。

调优点 4：预编译与 JIT

DeepGEMM 默认在第一次调用时 JIT 编译每个 GEMM 形状的 kernel——首次调用会有几百 ms 的 latency spike。生产部署前应做一遍”warmup”——用真实形状调用一次，让 JIT 缓存生成。否则第一个用户请求会被这个 spike 影响。

这四个调优点是部署 V4 时最常被忽略的”魔鬼细节”——但在大规模生产中累积起来差异显著。

13.9·延展 DeepGEMM 与 vLLM 中现有 GEMM 路径的并存

vLLM 已经有自己的 GEMM 实现——通常用 cuBLAS 或 CUTLASS。V4 集成 DeepGEMM 时不能简单”全部替换”——必须让两套 GEMM 路径并存。

为什么不能全部替换：vLLM 支持多个模型（Llama、Mistral、Qwen 等），它们用 BF16 / FP16 GEMM。如果把 cuBLAS 全部换成 DeepGEMM，这些模型反而变慢——DeepGEMM 没有针对 BF16 dense 模型优化。

vLLM 的策略：在 attention backend / MoE 模块里根据模型 dtype 动态 dispatch：

• 如果 weight.dtype == FP4_e2m1：走 DeepGEMM 的 fp4_gemm
• 如果 weight.dtype == FP8_e4m3：走 DeepGEMM 的 fp8_gemm
• 否则走 cuBLAS / CUTLASS

这种 dispatch 让 V4 部署不影响其他模型——一个 vLLM 实例可以同时跑 V4 + Llama 而不冲突。

编译期 dispatch 还是运行期？：dispatch 必须在运行期——因为同一个 vLLM 实例可能会动态加载不同模型。这意味着每次 GEMM 调用前要做 dtype 检查——开销在纳秒级，可以忽略。

与 PagedAttention 的协调：DeepGEMM 处理 attention 内部的 Q/K/V/O linear，但 PagedAttention 处理 attention 计算本身。两者完全解耦——一个负责”线性投影”、一个负责”注意力计算”。集成时不需要让它们共享代码。

第 19 章已经讲了部署的工程接缝——这里再次强调 DeepGEMM 与 vLLM 现有 GEMM 的并存——这是部署 V4 后其他模型也要继续跑 的工程现实。

13.9·拓展 DeepGEMM 的”开发者工效” 哲学

读完 DeepGEMM 仓库后，会注意到它的另一个工程哲学——对开发者的工效极致优化。具体体现：

工效点 1：每个 .cu 文件可独立读懂

DeepGEMM 不像 cutlass 那样模板嵌套——每个 .cu 文件可以独立打开、独立读完。一个工程师上手一周可以读完核心 600 行，理解每个 kernel 的设计。这种”可上手性”对开源项目极重要。

工效点 2：测试驱动开发

DeepGEMM 的 tests/ 目录覆盖每种 GEMM 形状、每种精度组合。任何对 kernel 的修改都先通过 tests——避免”这次改快了，下次改回退”的迭代灾难。

工效点 3：性能基准与正确性测试分离

tests/ 测正确性，benchmark/ 测性能。这种分离让”性能调优”与”功能开发”各自独立——调优时不需要担心改坏功能、开发时不需要担心降低性能。

工效点 4：与上游模型代码紧密协同

DeepGEMM 的接口（fp8_gemm 等）直接与 V4 的 inference/model.py 对接。这种”工具链一致性”让 V4 团队迭代时可以同时改模型代码 + GEMM 内核，不会遇到”两边接口不匹配”的工程债务。

工效点 5：编译时间快

DeepGEMM 的 CUDA 编译通常在几分钟内完成（单 GPU 架构）。对比 cutlass 的几十分钟编译时间，DeepGEMM 让”修改代码 → 测试” 的开发循环极短——工程师生产力高 5-10 倍。

这 5 个工效点是 DeepGEMM 在 V3 时代积累、V4 时代成熟的工程美学。它们让 DeepSeek 团队能在每代模型发布时同步发布配套 GEMM 库——不会出现”模型先发布、GEMM 三个月后才跟上” 的尴尬。

13.10 延伸阅读

• DeepGEMM 仓库 README：本章主要参考
• cutlass 文档：理解 GEMM kernel 的通用框架
• NVIDIA Hopper Architecture Whitepaper：WGMMA / TMA 指令细节
• 本书第 12 章：FP4 / FP8 / ue8m0 的格式细节
• 本书第 14 章：QAT 训练时的 act_quant 与 GEMM 配合

13.10·补 DeepGEMM 在”未来硬件”上的演进路径

DeepGEMM 当前覆盖 SM90 / SM100，但 NVIDIA 还有更多硬件路线图。把 DeepGEMM 在未来硬件上的演进做个推测。

B300（推测 2026 年下半年）：

下一代 NVIDIA 数据中心 GPU。预期改进：

• TensorCore 对 FP4 / FP6 的进一步原生加速
• 更大 SMEM（256-512 KB / SM）
• 改进的 TMA 指令

DeepGEMM 在 B300 上需要新增 sm_110 路径——预期工作量类似从 SM90 到 SM100 的增量。如果 V5 与 B300 同期发布，DeepGEMM 大概率在 V5 release 时同步更新。

B100（推测 2027 年）：

更高端的训练 GPU。预期改进：

• 更高 NVLink 带宽（800+ GB/s）
• 更大显存（300+ GB）

DeepGEMM 在 B100 上的主要工作是”调优新的 tile size 与 SMEM 配置”——算法不变。

国产 GPU（如华为昇腾 / 寒武纪）：

V4 的 README 提到”close integration with Huawei chips”——意味着 DeepGEMM 可能会有华为昇腾的版本（比如基于 CANN 的实现）。这部分代码可能不在公开 DeepGEMM 仓库——而是华为内部维护。

消费级 GPU（5090 / 6090）：

消费级 GPU 没有 NVLink + 显存有限——DeepGEMM 在这上面的优化优先级低。社区可能会有”消费级版本”的 DeepGEMM fork，但不会是 V4 团队维护。

理解这些演进路径让你做”硬件投资规划”——如果你 2026 下半年要建 V4 集群，B200 是合适的；如果是 2027 年，B300 / B100 可能更划算。

13.10·补·补 DeepGEMM 工程师速记

版本与依赖：

• CUDA 12.8+
• GCC 11+
• 必须 H100/H800/B200，不支持 A100
• 编译 5-10 分钟（首次）+ JIT 编译每个新 GEMM 形状（首次调用）

核心 API：

• fp8_gemm(x, x_scale, w, w_scale, scale_dtype)：FP8 e4m3 GEMM
• fp4_gemm(...)：FP4 e2m1 GEMM
• grouped_gemm_fp8(...)：MoE 用的 grouped GEMM
• act_quant(x, block_size, scale_fmt, scale_dtype, in_place)：激活量化

性能基线（M=N=K=4096）：

• H100：1300 TFlops（FP8）
• B200：~2000 TFlops（FP8 / FP4）

与 vLLM / SGLang 的集成：

• 自动按 weight.dtype dispatch
• FP4 / FP8 走 DeepGEMM、BF16 走 cuBLAS / cutlass
• 不需要用户改业务代码

调试工具：

• python benchmark/bench_fp8_gemm.py 跑性能基准
• pytest tests/ 跑正确性测试
• nsys profile 看 kernel 时间分布

常见问题：

• 编译失败：检查 CUDA / GCC 版本
• 性能不到峰值：检查 GPU 时钟、温度、Stream 配置
• 输出错误：检查 weight.scale 是否正确挂到 weight 张量上

13.10·延展 DeepGEMM 的”异步流水线” 优化

DeepGEMM 接近硬件极限的关键是异步流水线——把数据加载与计算重叠，让 GPU 不空闲。把这条流水线展开。

传统同步 GEMM 的问题：

Cycle 1: 从 global memory 加载 tile 0 到 SMEM
Cycle 2: GPU 等待加载完成
Cycle 3: TensorCore 计算 tile 0
Cycle 4: 加载 tile 1
Cycle 5: 等待
Cycle 6: 计算 tile 1

约一半的 cycle 在等内存——GPU TensorCore 利用率 50%。

DeepGEMM 的异步流水线：

Cycle 1: 启动 cp.async.bulk 加载 tile 0
Cycle 2: 启动加载 tile 1, 同时 wait tile 0
Cycle 3: 计算 tile 0, 同时启动加载 tile 2, wait tile 1
Cycle 4: 计算 tile 1, 同时启动加载 tile 3, wait tile 2
...

加载与计算重叠——TensorCore 持续工作，利用率 90%+。

实现关键：

• cp.async.bulk：H100 / B200 的异步拷贝指令，不阻塞 SM
• TMA descriptor：预定义的拷贝模板，启动开销极低
• multi-stage SMEM：SMEM 切成多份，正在用的 + 正在加载的 + 备用
• fence + wait：精确控制同步点，避免 race

与 sparse_attn 的串行依赖：

DeepGEMM 处理 attention 的 Q/K/V/O linear，sparse_attn 处理 attention 计算本身。两者是串行的——linear 必须先完成才能开始 attention。这个串行点是 V4 单层 layer 的延迟下限。

优化方向：

理论上 linear 与 attention 可以做”算子融合”——把 Q linear 与 attention 的 Q 投影合并到一个 kernel。FlashAttention v3 已经做了部分融合。DeepGEMM 在这方向还有空间——是 V5 / 未来 DeepGEMM 的优化方向。

理解流水线让你 debug 性能问题——如果某个 GEMM 跑得慢，可能是流水线被打断（如 SMEM 不够、依赖未满足）。

13.11 本章小结

• DeepGEMM 是为 V3/V4 量身定制的 FP8 / FP4 GEMM 库——cuBLAS 不能替代
• 关键设计：per-block 128×128 + ue8m0 scale + grouped GEMM + 与 act_quant 融合
• H100 路径用 WGMMA 指令做 FP8 GEMM、FP4 走”反量化 → FP8 GEMM”模拟
• B200 路径有原生 FP4 TensorCore 指令——比 H100 快约 1.6 倍
• 在 H100 上 FP8 GEMM 吞吐约 1300 TFlops，比 cuBLAS 快 30-40%
• 集成到 vLLM 等推理引擎是 V4 部署的必经之路

第 14 章：QAT 与 act_quant——V4 训练时的”假量化”全链路。