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 FP8W8A8 分块量化稠密 GEMM 实现的性能与精度DeepGEMMDeepSeek 开源的 DeepGEMM 库中的 block FP8 内核通过 vLLM 的兼容包装层 vllm/utils/deep_gemm.py 调用vLLM TritonvLLM 自研的 Triton 实现w8a8_triton_block_scaled_mm位于 fp8_utils.pyvLLM CUTLASSC/CUDA 扩展中的ops.cutlass_scaled_mm路径。原 README 明确界定了范围目前仅覆盖稠密denseGEMM且仅在 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 实测。二、环境准备SetupREADME 给出的安装步骤是按常规方式安装 vLLM 后单独从源码安装 DeepGEMMgit 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_GEMM1的默认环境之上。三、运行方式与输出解读Usage直接执行基准脚本python benchmark_fp8_block_dense_gemm.pyREADME 记录了 H10080GB 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/N4096,K7168,device_nameNVIDIA_H100_80GB_HBM3,dtypefp8_w8a8,block_shape[128,128].json INFO ... [fp8_utils.py:449] Using configuration from .../configs/N7168,K18432,device_nameNVIDIA_H100_80GB_HBM3,dtypefp8_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_M64、GROUP_SIZE_M32、num_warps4、num_stages2见 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-L278README 中贴出的输出对应的是较早一版形状清单含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), devicecuda, dtypetorch.bfloat16) B torch.randn((n, k), devicecuda, dtypetorch.bfloat16) C_ref A B.t() # BF16 参考结果 B_deepgemm, B_scale_deepgemm per_block_cast_to_fp8(B, [128, 128], use_ue8m0True) A_deepgemm, A_scale_deepgemm per_token_group_quant_fp8( A, block_size[1], column_major_scalesTrue, tma_aligned_scalesTrue) 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_scalesTrue)关键点在于三种实现对 scale比例因子张量的内存布局要求不同实现权重 B 量化激活 A 量化输出DeepGEMMper_block_cast_to_fp8(B, [128,128], use_ue8m0True)行主 列主 scale TMA 对齐写入预分配C_deepgemmvLLM Triton同 DeepGEMM 的 B可复用默认行主布局新张量vLLM CUTLASS同左列主 scale非 TMA 对齐新张量权重侧per_block_cast_to_fp8vllm/utils/deep_gemm.py#L719-L740把权重按 128×128 分块每块取amax / fp8_max作为 scaleuse_ue8m0True时先将 scale向上取整到 2 的幂UE8M0 格式这是 DeepGEMM 在 Hopper/Blackwell 上走硬件友好的比例因子路径所要求的。激活侧per_token_group_quant_fp8fp8_utils.py#L554-L680按group_size128做 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_dtypetorch.bfloat16, ) def vllm_cutlass_gemm(): return ops.cutlass_scaled_mm( A_vllm_cutlass, B_vllm.T, scale_aA_scale_vllm_cutlass, scale_bB_scale_vllm.T, out_dtypetorch.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 等并在支持 PDLSM90的架构上自动启用。Triton 路径w8a8_triton_block_scaled_mmfp8_utils.py#L880-L1008启动 Triton 内核_w8a8_triton_block_scaled_mmL756-L838。其核心模式是以BLOCK_SIZE_K沿 K 维循环每次迭代执行tl.dot(a, b)后立即乘以对应的 A/B scaleaccumulator 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.Tscale 以转置形式传入。这三条路径也正是 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 计时与指标口径计时循环采用warmup100次预热 repeat10000次计时迭代前后夹torch.accelerator.synchronize()指标计算如下benchmark_fp8_block_dense_gemm.py#L135-L139TFLOPS2·m·n·k / t / 1e12即一次 GEMM 的浮点运算量除以耗时GB/s(m·k k·n m·n·2) / t把 FP8 输入各 1 字节/元素与 BF16 输出2 字节/元素的总搬运量除以耗时作为有效内存带宽参考。对于小m的 decode 型形状如 m8、n/k 很大GB/s 往往比 TFLOPS 更能反映瓶颈所在——此时内核接近纯带宽受限。4.4 正确性度量为什么不是 assert_close精度表使用的calc_diff并非逐元素误差而是一个余弦相似性风格的全局指标vllm/utils/deep_gemm.py#L743-L756x, 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 应用时机量级上远小于量化本身的误差。另外注意脚本开头显式开启了 TF32torch.backends.cuda.matmul.allow_tf32 True并固定随机种子 42前者保证 BF16 参考矩阵乘的口径一致后者保证输入可复现。五、结论如何读综合 README 的原始数据与本节源码分析可以提炼出几条有实操价值的判断均以该次 H100 历史运行为据复现时请以本地数字为准大 mprefill场景DeepGEMM 在 m≥128 的大形状上稳定领先 Triton 2~3 倍以上如 m128、n18432、k7168 时 0.30x即 Triton 慢约 3.3 倍CUTLASS 与 DeepGEMM 互有胜负差距在 0.77x~0.90x 之间。小 m、瘦 n/k如 n24576/k1536、n32768/k512 的投影形状DeepGEMM 反而明显慢于 Triton 和 CUTLASS0.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.pyDeepGEMM 兼容包装fp8_gemm_nt / calc_diff / per_block_cast_to_fp8 / 支持性判断vllm/utils/deep_gemm.pyTriton 块量化 GEMM per-token-group 量化 JSON 配置查找vllm/model_executor/layers/quantization/utils/fp8_utils.pyTriton 内核调优配置目录vllm/model_executor/layers/quantization/utils/configs/线性层 DeepGEMM 内核 dispatchvllm/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），仅供参考