vLLM DeepGEMM FP8 Block GEMM 基准测试：DeepSeek DeepGEMM 对比 Triton 与 CUTLASS 的块量化 GEMM 内核解析

2026-09-06 20:42:03作者：龚格成

本文以 benchmarks/kernels/deepgemm/README.md 及其配套的 benchmark_fp8_block_dense_gemm.py 为核心，完整讲解如何在 vLLM 仓库中搭建环境、运行 DeepSeek DeepGEMM 与 vLLM 自带 Triton / CUTLASS 块量化 FP8 稠密 GEMM 内核的基准对比，并结合 vllm/utils/deep_gemm.py、fp8_utils.py 等源码剖析三条内核路径的量化布局差异、正确性度量方式与结果解读方法。读完本文，你能够独立复现该基准、读懂输出中的每张表格，并理解 vLLM 在生产路径上是如何在三种 FP8 GEMM 后端之间做选择的。

一、基准测试的对象与范围

该目录的目标是：在同一组矩阵形状下，横向对比三种 block FP8（W8A8 分块量化）稠密 GEMM 实现的性能与精度：

DeepGEMM：DeepSeek 开源的 DeepGEMM 库中的 block FP8 内核，通过 vLLM 的兼容包装层 vllm/utils/deep_gemm.py 调用；
vLLM Triton：vLLM 自研的 Triton 实现 w8a8_triton_block_scaled_mm，位于 fp8_utils.py；
vLLM CUTLASS：C++/CUDA 扩展中的 ops.cutlass_scaled_mm 路径。

原 README 明确界定了范围：目前仅覆盖稠密（dense）GEMM，且仅在 Hopper 架构 GPU 上有效。这是一个重要的适用前提——块量化 FP8 GEMM 是 DeepSeek 系列模型 FP8 推理的核心算子，128×128 分块量化（weight 按块、activation 按 token 分组）决定了三种实现在内存带宽与算力之间的平衡方式各不相同。

从当前源码结构看，范围实际上已在演进：vllm/utils/deep_gemm.py#L109-L115 中 is_deep_gemm_supported() 的注释写明“Currently, only Hopper and Blackwell GPUs are supported”，即当前 vLLM 的 DeepGEMM 包装层同时支持 Hopper 与 Blackwell 两代架构，而本 README 记录的是 H100 上的稠密 GEMM 实测。

二、环境准备（Setup）

README 给出的安装步骤是：按常规方式安装 vLLM 后，单独从源码安装 DeepGEMM：

git clone --recursive https://github.com/deepseek-ai/DeepGEMM
cd DeepGEMM
python setup.py install
uv pip install -e .

其中 --recursive 用于拉取 DeepGEMM 的子模块（其内核源码依赖子模块），安装需在 DeepGEMM 自己的目录内完成，避免污染 vLLM 的环境。

两个值得了解的细节，来自源码：

导入优先级：vllm/utils/deep_gemm.py#L175-L208 的 _import_deep_gemm() 会优先使用外部 pip 安装的 deep_gemm 包，找不到时才回退到 vLLM wheel 内自带的 vendored 副本 vllm.third_party.deep_gemm。这正是 README 要求单独安装 DeepGEMM 的原因——你可以固定（pin）一个特定的 DeepGEMM 版本，并覆盖 vLLM 内置版本。
JIT 缓存目录：_lazy_init() 在未设置 DG_JIT_CACHE_DIR 时，会将其指向 VLLM_CACHE_ROOT/deep_gemm（见 vllm/utils/deep_gemm.py#L265-L270）。DeepGEMM 内核是 JIT 编译的，首次运行会触发编译，后续复用缓存。
相关开关：vllm/envs.py 定义了 VLLM_USE_DEEP_GEMM（默认 1）、VLLM_USE_DEEP_GEMM_E8M0（默认 1，控制是否使用 UE8M0 幂次比例因子格式）和 VLLM_USE_DEEP_GEMM_TMA_ALIGNED_SCALES（默认 1）。本基准脚本正是构建在 VLLM_USE_DEEP_GEMM=1 的默认环境之上。

三、运行方式与输出解读（Usage）

直接执行基准脚本：

python benchmark_fp8_block_dense_gemm.py

README 记录了 H100（80GB HBM3）上的一次完整运行输出。完整输出分为五个部分，下面按出现顺序逐一拆解，并附上 README 中的原始数据。

3.1 环境与内核配置日志

输出的开头打印 PyTorch / CUDA / Triton 版本与设备名，随后出现若干关于 W8A8 Block FP8 内核配置的 WARNING/INFO，例如：

WARNING ... [fp8_utils.py:458] Using default W8A8 Block FP8 kernel config. Performance might be sub-optimal! Config file not found at .../configs/N=4096,K=7168,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
INFO ... [fp8_utils.py:449] Using configuration from .../configs/N=7168,K=18432,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json for W8A8 Block FP8 kernel.

这些日志直接对应 Triton 路径的配置查找逻辑 get_w8a8_block_fp8_configs：它按 N、K、设备名、block_shape 拼出 JSON 文件名，在 vllm/model_executor/layers/quantization/utils/configs/ 目录中查找针对该形状调优过的 Triton 内核参数（BLOCK_SIZE_M/N/K、num_warps 等）；找到就按最优配置运行，找不到则打印 WARNING 并回退到默认配置（默认 BLOCK_SIZE_M=64、GROUP_SIZE_M=32、num_warps=4、num_stages=2，见 fp8_utils.py#L962-L977）。因此解读 Triton 列的数字时需注意：日志里出现 “Config file not found” 的行，说明该形状用的是未调优的默认配置，Triton 侧性能可能被低估。

3.2 逐形状性能表（README 原始数据）

DeepGEMM Implementation：

+------+-------+-------+-----------+--------+--------+
| m    | n     | k     | Time (μs) | TFLOPS | GB/s   |
+------+-------+-------+-----------+--------+--------+
|    8 |  4096 |  7168 | 102.9     | 4.6    | 286.4  |
|    8 |  7168 | 18432 | 70.8      | 29.8   | 1868.8 |
|    8 | 18432 |  7168 | 69.3      | 30.5   | 1911.8 |
|   64 |  4096 |  7168 | 69.1      | 54.4   | 439.0  |
|   64 |  7168 | 18432 | 69.4      | 243.6  | 1933.6 |
|   64 | 18432 |  7168 | 70.4      | 240.3  | 1917.2 |
|   64 | 24576 |  1536 | 70.1      | 68.9   | 584.6  |
|   64 | 32768 |   512 | 68.4      | 31.4   | 307.1  |
|   64 |  7168 | 16384 | 69.5      | 216.3  | 1718.5 |
|  128 |  4096 |  7168 | 141.1     | 53.3   | 222.1  |
|  128 |  7168 | 18432 | 71.9      | 470.5  | 1896.1 |
|  128 | 18432 |  7168 | 69.3      | 488.2  | 1988.2 |
| 1024 |  4096 |  7168 | 89.7      | 670.1  | 502.5  |
| 1024 | 18432 |  7168 | 279.0     | 969.8  | 635.2  |
| 2048 |  4096 |  7168 | 175.1     | 687.0  | 347.4  |
| 4096 |  4096 |  7168 | 335.4     | 717.0  | 275.1  |
+------+-------+-------+-----------+--------+--------+

vLLM Triton Implementation：

+------+-------+-------+-----------+--------+--------+--------------+
| m    | n     | k     | Time (μs) | TFLOPS | GB/s   | vs DeepGEMM  |
+------+-------+-------+-----------+--------+--------+--------------+
|    8 |  4096 |  7168 | 74.0      | 6.3    | 398.2  | 1.39x faster |
|    8 |  7168 | 18432 | 89.6      | 23.6   | 1478.1 | 0.79x slower |
|    8 | 18432 |  7168 | 113.2     | 18.7   | 1170.4 | 0.61x slower |
|   64 |  4096 |  7168 | 79.4      | 47.3   | 382.2  | 0.87x slower |
|   64 |  7168 | 18432 | 98.5      | 171.7  | 1363.0 | 0.70x slower |
|   64 | 18432 |  7168 | 119.5     | 141.5  | 1129.4 | 0.59x slower |
|   64 | 24576 |  1536 | 37.6      | 128.4  | 1089.7 | 1.86x faster |
|   64 | 32768 |   512 | 38.7      | 55.5   | 542.6  | 1.77x faster |
|   64 |  7168 | 16384 | 86.1      | 174.5  | 1386.4 | 0.81x slower |
|  128 |  4096 |  7168 | 90.7      | 82.9   | 345.4  | 1.56x faster |
|  128 |  7168 | 18432 | 144.0     | 234.9  | 946.9  | 0.50x slower |
|  128 | 18432 |  7168 | 229.5     | 147.4  | 600.1  | 0.30x slower |
| 1024 |  4096 |  7168 | 242.3     | 248.2  | 186.1  | 0.37x slower |
| 1024 | 18432 |  7168 | 897.8     | 301.4  | 197.4  | 0.31x slower |
| 2048 |  4096 |  7168 | 463.0     | 259.7  | 131.4  | 0.38x slower |
| 4096 |  4096 |  7168 | 901.8     | 266.7  | 102.3  | 0.37x slower |
+------+-------+-------+-----------+--------+--------+--------------+

vLLM CUTLASS Implementation：

+------+-------+-------+-----------+--------+--------+--------------+--------------+
| m    | n     | k     | Time (μs) | TFLOPS | GB/s   | vs DeepGEMM  | vs Triton    |
+------+-------+-------+-----------+--------+--------+--------------+--------------+
|    8 |  4096 |  7168 | 34.6      | 13.6   | 852.3  | 2.98x faster | 2.14x faster |
|    8 |  7168 | 18432 | 78.9      | 26.8   | 1677.3 | 0.90x slower | 1.13x faster |
|    8 | 18432 |  7168 | 81.2      | 26.0   | 1631.1 | 0.85x slower | 1.39x faster |
|   64 |  4096 |  7168 | 36.9      | 101.9  | 822.9  | 1.87x faster | 2.15x faster |
|   64 |  7168 | 18432 | 87.4      | 193.4  | 1535.2 | 0.79x slower | 1.13x faster |
|   64 | 18432 |  7168 | 85.0      | 199.0  | 1587.6 | 0.83x slower | 1.41x faster |
|   64 | 24576 |  1536 | 28.0      | 172.8  | 1465.8 | 2.51x faster | 1.35x faster |
|   64 | 32768 |   512 | 28.8      | 74.5   | 728.5  | 2.37x faster | 1.34x faster |
|   64 |  7168 | 16384 | 77.9      | 193.0  | 1532.8 | 0.89x slower | 1.11x faster |
|  128 |  4096 |  7168 | 39.1      | 192.4  | 802.0  | 3.61x faster | 2.32x faster |
|  128 |  7168 | 18432 | 93.7      | 360.8  | 1454.2 | 0.77x slower | 1.54x faster |
|  128 | 18432 |  7168 | 85.7      | 394.8  | 1608.0 | 0.81x slower | 2.68x faster |
| 1024 |  4096 |  7168 | 99.7      | 603.1  | 452.2  | 0.90x slower | 2.43x faster |
| 1024 | 18432 |  7168 | 331.3     | 816.7  | 534.9  | 0.84x slower | 2.71x faster |
| 2048 |  4096 |  7168 | 198.3     | 606.6  | 306.7  | 0.88x slower | 2.34x faster |
| 4096 |  4096 |  7168 | 392.2     | 613.2  | 235.3  | 0.86x slower | 2.30x faster |
+------+-------+-------+-----------+--------+--------+--------------+--------------+

形状选择值得留意：m ∈ {8, 64, 128} 模拟解码（decode）阶段的小批量场景，m ∈ {1024, 2048, 4096} 模拟预填充（prefill）阶段的大批量场景；n 与 k 取自真实 LLM 层的投影维度（如 4096/7168/18432/24576/32768 等）。从源码结构看，脚本当前实际生效的形状清单是第二处 shapes 赋值（16 个形状，m ∈ {64, 128, 4096}，见 benchmark_fp8_block_dense_gemm.py#L259-L278），README 中贴出的输出对应的是较早一版形状清单（含 m = 8 / 1024 / 2048）。复现时实际打印的形状以脚本当前定义为准。

3.3 平均性能、平均加速比与精度

README 输出的最后三部分：

===== AVERAGE PERFORMANCE =====
+----------------+------------+----------+---------------+
| Implementation | Avg TFLOPS | Avg GB/s | Avg Time (ms) |
+----------------+------------+----------+---------------+
| DeepGEMM       | 310.98     | 1052.10  | 0.11          |
| vLLM Triton    | 144.30     | 715.60   | 0.23          |
| vLLM CUTLASS   | 286.78     | 1076.67  | 0.11          |
+----------------+------------+----------+---------------+

===== AVERAGE SPEEDUPS =====
+-----------------------------+--------------+
| Comparison                  | Speedup      |
+-----------------------------+--------------+
| DeepGEMM vs vLLM Triton     | 1.71x faster |
| DeepGEMM vs vLLM CUTLASS    | 0.94x slower |
| vLLM CUTLASS vs vLLM Triton | 1.84x faster |
+-----------------------------+--------------+

===== ACCURACY COMPARISON =====
+----------------+-----------------------+
| Implementation | Avg Diff vs Reference |
+----------------+-----------------------+
| DeepGEMM       | 0.000684              |
| vLLM Triton    | 0.000684              |
| vLLM CUTLASS   | 0.000684              |
+----------------+-----------------------+

这组数据的含义是：在该次 H100 实测中，DeepGEMM 平均比 vLLM Triton 快 1.71 倍；与 vLLM CUTLASS 相比略慢（0.94x）；而 CUTLASS 平均比 Triton 快 1.84 倍；三者相对 BF16 参考结果的平均偏差完全一致（0.000684），说明精度上没有区分度。这些是特定历史运行的快照数据，实际数值会随 DeepGEMM 版本、vLLM 内核更新与具体卡型而变化，请以本地复现结果为准。

四、源码级剖析：三条内核路径到底测了什么

基准脚本 benchmark_fp8_block_dense_gemm.py 的核心是 benchmark_shape()（L22-L183）。它保证了对比公平性：三个实现消费同一对 BF16 输入矩阵，且量化步骤在计时之外预先完成（权重 B 可离线预量化，激活 A 的量化则按各内核要求生成不同布局）。

4.1 参考值与量化预处理

A = torch.randn((m, k), device="cuda", dtype=torch.bfloat16)
B = torch.randn((n, k), device="cuda", dtype=torch.bfloat16)
C_ref = A @ B.t()                                    # BF16 参考结果
B_deepgemm, B_scale_deepgemm = per_block_cast_to_fp8(B, [128, 128], use_ue8m0=True)
A_deepgemm, A_scale_deepgemm = per_token_group_quant_fp8(
    A, block_size[1], column_major_scales=True, tma_aligned_scales=True)
A_vllm, A_scale_vllm = per_token_group_quant_fp8(A, block_size[1])
A_vllm_cutlass, A_scale_vllm_cutlass = per_token_group_quant_fp8(
    A, block_size[1], column_major_scales=True)

关键点在于三种实现对 scale（比例因子）张量的内存布局要求不同：

实现	权重 B 量化	激活 A 量化	输出
DeepGEMM	`per_block_cast_to_fp8(B, [128,128], use_ue8m0=True)`	行主 + 列主 scale + TMA 对齐	写入预分配 `C_deepgemm`
vLLM Triton	同 DeepGEMM 的 B（可复用）	默认行主布局	新张量
vLLM CUTLASS	同左	列主 scale（非 TMA 对齐）	新张量

权重侧：per_block_cast_to_fp8（vllm/utils/deep_gemm.py#L719-L740）把权重按 128×128 分块，每块取 amax / fp8_max 作为 scale；use_ue8m0=True 时先将 scale 向上取整到 2 的幂（UE8M0 格式），这是 DeepGEMM 在 Hopper/Blackwell 上走硬件友好的比例因子路径所要求的。
激活侧：per_token_group_quant_fp8（fp8_utils.py#L554-L680）按 group_size=128 做 per-token-group 量化，在 CUDA 平台上优先调用 _C.per_token_group_fp8_quant 原生内核，否则回退 Triton 内核；column_major_scales 与 tma_aligned_scales 参数专门满足 DeepGEMM 对 scale 张量“列主 + TMA 对齐”的布局约束。

4.2 三个被测函数

def deepgemm_gemm():
    fp8_gemm_nt(
        (A_deepgemm, A_scale_deepgemm), (B_deepgemm, B_scale_deepgemm), C_deepgemm
    )
    return C_deepgemm

def vllm_triton_gemm():
    return w8a8_triton_block_scaled_mm(
        A_vllm, B_vllm, A_scale_vllm, B_scale_vllm,
        block_size, output_dtype=torch.bfloat16,
    )

def vllm_cutlass_gemm():
    return ops.cutlass_scaled_mm(
        A_vllm_cutlass, B_vllm.T,
        scale_a=A_scale_vllm_cutlass, scale_b=B_scale_vllm.T,
        out_dtype=torch.bfloat16,
    )

DeepGEMM 路径：脚本导入的 fp8_gemm_nt 是 vLLM 的兼容包装（vllm/utils/deep_gemm.py#L460-L469），内部把 is_deep_gemm_e8m0_used() 的决策转换为底层 disable_ue8m0_cast 参数。包装层通过 _lazy_init() 延迟解析 DeepGEMM 的所有符号（dense GEMM、grouped GEMM、MQA logits 等），并在支持 PDL（SM90+）的架构上自动启用。
Triton 路径：w8a8_triton_block_scaled_mm（fp8_utils.py#L880-L1008）启动 Triton 内核 _w8a8_triton_block_scaled_mm（L756-L838）。其核心模式是：以 BLOCK_SIZE_K 沿 K 维循环，每次迭代执行 tl.dot(a, b) 后立即乘以对应的 A/B scale（accumulator += tl.dot(a, b) * a_s[:, None] * b_s[None, :]），并用 GROUP_SIZE_M 分组重排 block 顺序以改善 L2 命中。内核启动前会按第 3.1 节所述查 JSON 调优配置。
CUTLASS 路径：ops.cutlass_scaled_mm 是 vLLM CUDA 扩展暴露的块量化 scaled matmul，输入 B 需要转置为列主视图（B_vllm.T），scale 以转置形式传入。

这三条路径也正是 vLLM 生产推理中的候选内核：线性层的 dispatch 逻辑位于 vllm/model_executor/kernels/linear/scaled_mm/deep_gemm.py 等文件，选择 DeepGEMM 前置条件由 should_use_deepgemm_for_fp8_linear 给出——输出 dtype 必须是 BF16，且权重 N 维是 64 的倍数、K 维是 128 的倍数（这也解释了为什么基准形状都整齐地取 64/128 的倍数）。

4.3 计时与指标口径

计时循环采用 warmup=100 次预热 + repeat=10000 次计时迭代，前后夹 torch.accelerator.synchronize()，指标计算如下（benchmark_fp8_block_dense_gemm.py#L135-L139）：

TFLOPS = 2·m·n·k / t / 1e12，即一次 GEMM 的浮点运算量除以耗时；
GB/s = (m·k + k·n + m·n·2) / t，把 FP8 输入（各 1 字节/元素）与 BF16 输出（2 字节/元素）的总搬运量除以耗时，作为有效内存带宽参考。

对于小 m 的 decode 型形状（如 m=8、n/k 很大），GB/s 往往比 TFLOPS 更能反映瓶颈所在——此时内核接近纯带宽受限。

4.4 正确性度量：为什么不是 assert_close

精度表使用的 calc_diff 并非逐元素误差，而是一个余弦相似性风格的全局指标（vllm/utils/deep_gemm.py#L743-L756）：

x, y = x.double(), y.double()
denominator = (x * x + y * y).sum()
sim = 2 * (x * y).sum() / denominator
return 1 - sim

源码注释解释了原因：DeepGEMM 内核在 Blackwell/B200 上存在可见的逐元素误差，会导致 torch.testing.assert_close 失败，因此改用全局相似度 1 - sim 来报告整体偏差。三个实现得到完全相同的 0.000684，是因为它们消费相同的量化输入，差异仅来自内核内部的累加顺序与 scale 应用时机，量级上远小于量化本身的误差。

另外注意脚本开头显式开启了 TF32（torch.backends.cuda.matmul.allow_tf32 = True）并固定随机种子 42，前者保证 BF16 参考矩阵乘的口径一致，后者保证输入可复现。

五、结论如何读

综合 README 的原始数据与本节源码分析，可以提炼出几条有实操价值的判断（均以该次 H100 历史运行为据，复现时请以本地数字为准）：

大 m（prefill）场景：DeepGEMM 在 m≥128 的大形状上稳定领先 Triton 2~3 倍以上（如 m=128、n=18432、k=7168 时 0.30x，即 Triton 慢约 3.3 倍），CUTLASS 与 DeepGEMM 互有胜负（差距在 0.77x~0.90x 之间）。
小 m、瘦 n/k（如 n=24576/k=1536、n=32768/k=512 的投影形状）：DeepGEMM 反而明显慢于 Triton 和 CUTLASS（0.59x~0.30x 量级的反向差距），说明三种内核在“带宽受限 + 小批量”这一区间的最优布局策略不同。
Triton 路径受调优配置影响显著：README 日志显示部分形状缺少对应的 JSON 调优配置（回退默认配置），这类形状下 Triton 的劣势可能被进一步放大；configs 目录按 N,K,device_name,dtype,block_shape 粒度管理配置，新增卡型时应关注日志中是否出现 “Config file not found” 告警。
精度无差异：三者相对参考值的全局偏差一致，选型可以完全由性能驱动。

六、相关文件索引

内容	路径
基准文档（本文主体）	benchmarks/kernels/deepgemm/README.md
基准脚本	benchmarks/kernels/deepgemm/benchmark_fp8_block_dense_gemm.py
DeepGEMM 兼容包装（fp8_gemm_nt / calc_diff / per_block_cast_to_fp8 / 支持性判断）	vllm/utils/deep_gemm.py
Triton 块量化 GEMM + per-token-group 量化 + JSON 配置查找	vllm/model_executor/layers/quantization/utils/fp8_utils.py
Triton 内核调优配置目录	vllm/model_executor/layers/quantization/utils/configs/
线性层 DeepGEMM 内核 dispatch	vllm/model_executor/kernels/linear/scaled_mm/deep_gemm.py
相关环境变量定义（VLLM_USE_DEEP_GEMM 等）	vllm/envs.py