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

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

【免费下载链接】vllmA high-throughput and memory-efficient inference and serving engine for LLMs项目地址: https://gitcode.com/GitHub_Trending/vl/vllm

本文以 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

【免费下载链接】vllmA high-throughput and memory-efficient inference and serving engine for LLMs项目地址: https://gitcode.com/GitHub_Trending/vl/vllm

创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

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