蚂蚁集团与 vLLM 社区开源量化推理内核库 Humming

蚂蚁集团与 vLLM ç¤¾åŒºå¼€æºé‡åŒ–æŽ¨ç†å† æ ¸åº“ Humming · AIHOT

vLLM 官方博客

·

Ant Group and the vLLM Community

2026-10-21 08:00

7 小时前

精选

AI 评分

67

蚂蚁集团与 vLLM ç¤¾åŒºå¼€æºé‡åŒ–æŽ¨ç†å† æ ¸åº“ Humming

Announcing Humming: Fast, Flexible Kernels for Quantized Inference

AI 导读

蚂蚁集团与 vLLM ç¤¾åŒºå¼€æºé€šç”¨é«˜æ€§èƒ½é‡åŒ–æŽ¨ç†å† æ ¸åº“ Humming，仓库为 github.com/vllm-project/humming。

推荐理由

åŽŸæ–‡ç»™å‡ºç»Ÿä¸€é‡åŒ–è¡¨ç¤ºåŠ JIT ç¼–è¯‘çš„å† æ ¸è®¾è®¡å’Œå¤šé¡¹åŸºå‡†æ•°å­—ï¼Œè¯»è€ å¯ä»¥æ®æ­¤è¯„ä¼°ä¸åŒ GPU 上的量化推理方案选择。

正文 · AI 翻译

蚂蚁集团与 vLLM 社区开源 Hummingï¼Œè¿™æ˜¯ä¸€ä¸ªé¢å‘é‡åŒ–æŽ¨ç†çš„é€šç”¨é«˜æ€§èƒ½å† æ ¸åº“ã€‚Humming å ·å¤‡ä»¥ä¸‹ç‰¹æ€§ï¼š

• 广泛的设备支持 ：支持从 Turing 起的 NVIDIA GPU 架构，以及平头哥镇岳 810E 等 CUDA å ¼å®¹è®¾å¤‡ã€‚

• å ¨é¢çš„é‡åŒ–æ”¯æŒ ï¼šç»Ÿä¸€æž¶æž„æ”¯æŒç¨ å¯†ä¸Ž MoE å·¥ä½œè´Ÿè½½ä¸‹æƒé‡å’Œæ¿€æ´»æ ¼å¼çš„å¤šç§ç»„åˆï¼Œè¦†ç›–å¹¿æ³›çš„æ•´æ•°ä¸Žæµ®ç‚¹ç±»åž‹ã€ä½å®½åŠç¼©æ”¾æ–¹æ¡ˆã€‚

• 强劲的执行性能 ：在许多代表性工作负载上显著优于 vLLM Marlin、FlashInfer 和 DeepGEMM 的现有实现。

仓库： github.com/vllm-project/humming

背景

éšç€æ¨¡åž‹å’ŒæŽ¨ç†å·¥ä½œè´Ÿè½½çš„å¢žé•¿ï¼Œå† å­˜å®¹é‡ã€å† å­˜å¸¦å®½å’Œè®¡ç®—æˆæœ¬æ—¥ç›Šåˆ¶çº¦éƒ¨ç½²ã€‚è¿™åœ¨ MoE 模型中尤为明显：每个 token ä» æ¿€æ´»ä¸€éƒ¨åˆ†ä¸“å®¶ï¼Œä½†å ¨éƒ¨æƒé‡ä»éœ€å­˜å‚¨ã€‚ä¾‹å¦‚ DeepSeek-V3 拥有 6710 亿参数，每个 token 激活 370 äº¿å‚æ•°ã€‚é‡åŒ–é€šè¿‡é™ä½Žæƒé‡å’Œæ¿€æ´»çš„ç²¾åº¦æ¥å‡å°‘å† å­˜å ç”¨å’Œæ•°æ®æ¬è¿ï¼ŒåŒæ—¶ä¹Ÿèƒ½æå‡è®¡ç®—æ•ˆçŽ‡ã€‚

不同模型和 GPU ä»£é™ å¯¹é‡åŒ–çš„éœ€æ±‚å„ä¸ç›¸åŒã€‚å¸¸è§é ç½®åŒ æ‹¬ WNA16、W8A8、W4A8 和 W4A4ï¼Œå ¶ä¸­ W 和 A 分别表示权重和激活位宽，N 表示可变的权重位宽。即使精度相同，INT4 和 NVFP4 ç­‰æ–¹æ¡ˆåœ¨æ•°å€¼æ ¼å¼å’Œç¼©æ”¾æ–¹å¼ä¸Šä¹Ÿæœ‰å·®å¼‚ã€‚ç¼©æ”¾ç²’åº¦ã€ç¼©æ”¾è¡¨ç¤ºã€åˆ†ç»„å¤§å°ä»¥åŠå¯¹ç§°æˆ–éžå¯¹ç§°é‡åŒ–è¿›ä¸€æ­¥æ‰©å±•äº†é ç½®ç©ºé—´ã€‚

åªæœ‰å€ŸåŠ©é«˜æ•ˆå† æ ¸ï¼Œè¿™äº›å­˜å‚¨å’Œè®¡ç®—ä¸Šçš„èŠ‚çœæ‰èƒ½è½¬åŒ–ä¸ºæ›´ä½Žçš„å»¶è¿Ÿå’Œæ›´é«˜çš„åžåã€‚Prefill、decode 以及不同的 MoE ä¸“å®¶è´Ÿè½½å¯¹è®¡ç®—å’Œå† å­˜æå‡ºä¸åŒè¦æ±‚ã€‚å† æ ¸å¿ é¡»é€‚åº”è¿™äº›å·¥ä½œè´Ÿè½½ï¼ŒåŒæ—¶ä¿æŒè½¬æ¢ã€åé‡åŒ–å’Œè°ƒåº¦å¼€é”€å¤„äºŽä½Žä½ã€‚

ä¸ºæ¯ç§ç»„åˆå¼€å‘å’Œé›†æˆä¸“ç”¨å† æ ¸ä¼šå¸¦æ¥ä¸æ–­å¢žé•¿çš„ç»´æŠ¤è´Ÿæ‹ ã€‚æ¯æ–°å¢žä¸€ç§ç»„åˆï¼Œéƒ½å¯èƒ½éœ€è¦å„è‡ªçš„æƒé‡é¢„å¤„ç†ã€å† æ ¸é€‰æ‹©ã€æ¡†æž¶é€‚é å™¨å’Œæµ‹è¯•ã€‚ä¼˜åŒ–å·¥ä½œéšåŽè¿˜éœ€åœ¨å¤šä¸ªå®žçŽ°ä¹‹é—´åŒæ­¥ã€‚

Humming é€šè¿‡å°†é‡åŒ–æè¿°ä¸Žæ‰§è¡Œé€‰æ‹©åˆ†ç¦»ï¼Œå¹¶ç”±å ±äº«ç»„ä»¶ç»„åˆå† æ ¸æ¥è§£å†³è¿™ä¸€é—®é¢˜ã€‚

通用架构

Humming çš„æ‰§è¡Œè·¯å¾„åˆ†ä¸ºä¸‰éƒ¨åˆ†ï¼šè¡¨ç¤ºé‡åŒ–è®¡ç®—ã€ä¸ºç¡¬ä»¶å’Œå·¥ä½œè´Ÿè½½é€‰æ‹©é ç½®ã€ä»Ž CUDA ç»„ä»¶ç¼–è¯‘å† æ ¸ã€‚

统一的量化表示

ä¸åŒé‡åŒ–å·¥å ·å¯¼å‡ºçš„æ¨¡åž‹åœ¨æ•°å€¼æ ¼å¼ã€å‚æ•°ç»„ç»‡å’Œæƒé‡å¸ƒå±€ä¸Šå„ä¸ç›¸åŒã€‚Humming å°†æ”¯æŒçš„å¯¼å‡ºæ ¼å¼è½¬æ¢ä¸º Humming æ ‡å‡†æ ¼å¼ï¼šä¸€ç§é€šç”¨çš„é‡åŒ–æè¿°ï¼Œæƒé‡ã€ç¼©æ”¾ã€é›¶ç‚¹åŠå ¶ä»–å¼ é‡å‡æŒ‰ä¸€è‡´æ–¹å¼ç»„ç»‡ã€‚

è¯¥è¡¨ç¤ºåœ¨ä¿ç•™å„æ–¹æ¡ˆé‡åŒ–è¯­ä¹‰çš„åŒæ—¶ï¼Œå°†æƒé‡æ ¼å¼ã€ç¼©æ”¾ç²’åº¦ã€ç¼©æ”¾æ•°æ®ç±»åž‹ã€é›¶ç‚¹åŠäºŒçº§ç¼©æ”¾æè¿°ä¸ºç‹¬ç«‹å‚æ•°ã€‚æ¿€æ´»æœ‰å ¶è‡ªèº«çš„æ•°å€¼æ ¼å¼å’Œé‡åŒ–è®¾ç½®ã€‚è¿™äº›å‚æ•°å¯åœ¨æ ¼å¼å’Œç¡¬ä»¶çº¦æŸä¸‹è¿›è¡Œç»„åˆã€‚

例如，分组 INT4 权重可以与 BF16 æ¿€æ´»å€¼é å¯¹ä½¿ç”¨ã€‚é‡‡ç”¨ä¸åŒç¼©æ”¾æ–¹æ¡ˆçš„ä½Žç²¾åº¦æµ®ç‚¹æƒé‡ä½¿ç”¨åŒä¸€å¥—å‚æ•°ä½“ç³»ã€‚è¿™ä¸ºæŽ¨ç†æ¡†æž¶æä¾›äº†é¢å‘ä¸åŒæ¥æºæ¨¡åž‹çš„ç»Ÿä¸€æŽ¥å£ã€‚

针对每种工作负载的调优与 JIT 编译

æ ‡å‡†æ ¼å¼åœ¨å„ç±»å¯¼å‡ºæ ¼å¼ä¹‹é—´æä¾›ç»Ÿä¸€çš„è¡¨ç¤ºã€‚éšåŽï¼Œé¢„å¤„ç†ä¼šæ ¡éªŒå‚æ•°ã€å¡«å ç»´åº¦ï¼Œå¹¶å°†è¿™äº›å¼ é‡é‡æ–°æ‰“åŒ ä¸ºç›®æ ‡å† æ ¸é€‚ç”¨çš„å¸ƒå±€ã€‚Humming 将量化描述与权重形状、GEMM ç±»åž‹ã€ç›®æ ‡ GPU å’Œè®¡ç®—é€‰é¡¹ç»“åˆèµ·æ¥ï¼Œä»¥é ç½®æ‰§è¡Œè¿‡ç¨‹ã€‚

ä¸‰ç§é ç½®ç”¨äºŽæè¿°å† æ ¸ï¼š

• å±‚é ç½® æè¿°å¼ é‡å½¢çŠ¶ã€æ•°æ®ç±»åž‹ã€é‡åŒ–å‚æ•°å’Œç›®æ ‡æž¶æž„ã€‚

• è®¡ç®—é ç½® 指定 GEMM ç±»åž‹å’Œè®¡ç®—é€‰é¡¹ï¼Œä¾‹å¦‚æ‰¹æ¬¡ä¸å˜æ€§å’Œè¾“å ¥ç¼©æ”¾å¸ƒå±€ã€‚

• è°ƒä¼˜é ç½® 选择瓦片形状、流水线阶段数、warp ç‰¹åŒ–ã€å† å­˜è®¿é—®å’Œ MMA 指令。它还控制并发与调度；TMA å¯ä»¥æŒ‰å¼ é‡è¿›è¡Œé ç½®ã€‚

é»˜è®¤æƒ å†µä¸‹ï¼ŒHumming ä½¿ç”¨é’ˆå¯¹å ·ä½“è®¾å¤‡çš„å¯å‘å¼è§„åˆ™æ¥é€‰æ‹©è°ƒä¼˜é ç½®ã€‚å®ƒå¯ä»¥ä¸ºä¸åŒèŒƒå›´çš„è¾“å ¥å°ºå¯¸å‡†å¤‡ä¸åŒçš„é ç½®ï¼Œå¹¶åœ¨è¿è¡Œæ—¶åˆ†æ´¾åˆ°åˆé€‚çš„å† æ ¸ã€‚

Humming 的统一架构使针对不同工作负载调优执行参数和选择流水线变得简单直接。在小批次尺寸下，使用异步拷贝（ cp.async ）和 MMA 的多级流水线可能优于围绕 TMA 和 UMMA 构建的流水线。在较大批次尺寸下，TMA 和 UMMA 可以提供更好的吞吐量。Humming åœ¨åŒä¸€é ç½®å’Œç»„ä»¶æ¡†æž¶å† åŒæ—¶æ”¯æŒè¿™ä¸¤ç§æ‰§è¡Œè·¯å¾„ï¼Œä»Žè€Œèƒ½å¤Ÿä¸ºæ¯ç§å·¥ä½œè´Ÿè½½é€‰æ‹©æ›´æœ‰æ•ˆçš„ç­–ç•¥ã€‚

è¿™ä¸‰ç§é ç½®å ±åŒç‰¹åŒ–ç”± NVRTC 或 NVCC ç¼–è¯‘çš„å† æ ¸ã€‚ç¼–è¯‘åŽçš„å† æ ¸ä¼šè¢«ç¼“å­˜ä»¥ä¾¿å¤ç”¨ã€‚æŽ¨ç†æ—¶å°†é€‰å®šçš„å† æ ¸ä¸Žé¢„å¤„ç†åŽçš„æƒé‡å’Œè¿è¡Œæ—¶æ¿€æ´»å€¼ç»“åˆèµ·æ¥ã€‚

è·¨æ ¼å¼ä¸Žæž¶æž„çš„å ±äº« CUDA 组件

在 CUDA 层面，Humming 将任务调度、数据移动、反量化与缩放、Tensor Core è®¡ç®—ä»¥åŠè¾“å‡ºå¤„ç†ç»„ç»‡ä¸ºå¯å¤ç”¨çš„ç»„ä»¶ã€‚å®ƒæä¾›æ ‡å‡†æµæ°´çº¿ã€warp 特化流水线和专用 UMMA æµæ°´çº¿ã€‚è¿™äº›æµæ°´çº¿å ±äº«å ¼å®¹çš„ç»„ä»¶ï¼ŒåŒæ—¶é‡‡ç”¨å„è‡ªæ‰§è¡Œç­–ç•¥æ‰€éœ€çš„æ•°æ®æµå’ŒåŒæ­¥æ–¹å¼ã€‚

每组组件处理不同的变化来源：

• æ•°å€¼å¤„ç†è´Ÿè´£é‡åŒ–æ ¼å¼ã€‚ è§£ç å’Œç®—æœ¯ç»„ä»¶å¤„ç†ä½å®½ã€æ•´æ•°æˆ–æµ®ç‚¹è¡¨ç¤ºã€ç¼©æ”¾å› å­å’Œé›¶ç‚¹ã€‚åé‡åŒ–ä¸Žç¼©æ”¾è¢«é›†æˆåˆ°å† æ ¸çš„è®¡ç®—ä¸­ã€‚

• å† å­˜ä¸Žè®¡ç®—åŽç«¯è´Ÿè´£ GPU 架构。 Humming ä¸ºç›®æ ‡æž¶æž„é€‰æ‹©æ•°æ®ç§»åŠ¨å’Œ Tensor Core 实现，并在支持时使用异步拷贝、TMA 和多播。

• ä»»åŠ¡æ˜ å°„ä¸Žè°ƒåº¦è´Ÿè´£ç¨ å¯†å’Œ MoE 工作负载。 每种 GEMM ç±»åž‹ä½¿ç”¨ç›¸åº”çš„è°ƒåº¦é€»è¾‘ã€‚ç¼©æ”¾å’Œåç½®åŠ æ³•ä¹Ÿå¯ä»¥èžåˆåˆ°è¾“å‡ºé˜¶æ®µã€‚

JIT ç¼–è¯‘ä¼šæ ¹æ®ä¸‰ç§é ç½®ç»„åˆè¿™äº›ç»„ä»¶ï¼Œå¹¶æ¶ˆé™¤æœªä½¿ç”¨çš„å®žçŽ°åˆ†æ”¯ã€‚å› æ­¤ï¼Œå¯ä»¥é€šè¿‡æ‰©å±•ç›¸å ³ç»„ä»¶æ¥æ”¯æŒæ–°çš„æ ¼å¼æˆ–æž¶æž„ã€‚å¯¹å ±äº«ç»„ä»¶çš„æ”¹è¿›ä¼šä½¿å ¼å®¹çš„é ç½®å—ç›Šï¼Œä»Žè€Œå‡å°‘å¯¹å„è‡ªä¸“ç”¨å† æ ¸çš„é‡å¤å¼€å‘å’Œç»´æŠ¤ã€‚

这个 CUDA å®žçŽ°è½»é‡ä¸”é¿å äº†é‡é‡çº§å† æ ¸æ¡†æž¶ä¾èµ–ï¼ŒåŒæ—¶ä¿æŒ JIT ç¼–è¯‘å¿«é€Ÿã€‚åœ¨æˆ‘ä»¬äºŽé å¤‡ CUDA 13.4 的 H200 服务器上进行的单 W4A16 å† æ ¸ç¼–è¯‘æµ‹è¯•ä¸­ï¼Œåœ¨ç¦ç”¨ç¼–è¯‘ç¼“å­˜çš„æƒ å†µä¸‹ï¼Œç¼–è¯‘æ—¶é—´ä»Ž CUTLASS é åˆ NVCC 的 23.35 秒 降至 Humming é åˆ NVCC 的 1.82 秒 或 é åˆ NVRTC 的 1.88 秒 ã€‚è¿™é™ä½Žäº†ä¸ºæ–°å·¥ä½œè´Ÿè½½å‡†å¤‡å† æ ¸çš„æˆæœ¬ï¼Œå¹¶ä½¿è¿­ä»£è°ƒä¼˜æ›´åŠ å®žç”¨ã€‚Humming è¿˜ä¼šåœ¨æ¨¡åž‹é¢„çƒ­æœŸé—´å¹¶è¡Œç¼–è¯‘å† æ ¸ã€‚åœ¨æˆ‘ä»¬çš„æµ‹è¯•ä¸­ï¼Œç¼–è¯‘å¹¶åŠ è½½è¦†ç›–å•å±‚ GEMM 所有受支持 M å¤§å°çš„å† æ ¸ä» éœ€ 约 2–4 秒 ï¼Œä½¿å¾—åœ¨æŽ¨ç†å‰å‡†å¤‡å¤šä¸ªä¸“ç”¨å† æ ¸å˜å¾—åˆ‡å®žå¯è¡Œã€‚

基于设备参数的性能调优

Humming 利用 GEMM é ç½®å’Œè®¾å¤‡å‚æ•°å¯¹å† æ ¸è¿›è¡Œæ·±åº¦ä¼˜åŒ–ã€‚é‡åŒ–æ ¼å¼ã€çŸ©é˜µå½¢çŠ¶ã€GPU 指令支持、SM æ•°é‡ã€å ±äº«å† å­˜å’Œå¯„å­˜å™¨èµ„æºéƒ½ä¼šå½±å“æ‰§è¡Œé€‰æ‹©ï¼Œä»Žè€Œå°†ä½Žä½å­˜å‚¨çš„ä¼˜åŠ¿è½¬åŒ–ä¸ºå®žé™ çš„æŽ¨ç†æ€§èƒ½ã€‚

在现有 GPU 上实现高效的 W4A8-MXFP4 推理。 Blackwell 之前的架构缺乏原生 MXFP4 Tensor Core æ”¯æŒï¼Œè€Œä¼ ç»Ÿçš„ W4A8-MXFP4 å† æ ¸åœ¨è¾ƒå¤§ M 时可能受限于反量化和缩放开销。Humming 通过快速的 MXFP4 到 INT8/FP8 åé‡åŒ–è·¯å¾„è§£å†³äº†è¿™ä¸€é—®é¢˜ï¼Œè¯¥è·¯å¾„åœ¨å† æ ¸å† å°† MXFP4 æƒé‡è½¬æ¢ä¸ºå—æ”¯æŒçš„è®¡ç®—æ ¼å¼ã€‚åœ¨æ­¤è®¾ç½®ä¸‹ï¼Œè¯¥è·¯å¾„åœ¨è¾ƒå° M æ—¶æ˜¾è‘—ä¼˜äºŽä¼ ç»Ÿ W4A8 实现，在较大 M 时接近 W8A8 æ€§èƒ½ã€‚å› æ­¤ï¼ŒH200 等 GPU 可以运行高性能 W4A8-MXFP4 å† æ ¸ï¼ŒåŒæ—¶ä¿ç•™ 4 位权重的存储优势。

ä¸ºæ¯ç§è®¾å¤‡åŒ¹é åé‡åŒ–è·¯å¾„ã€‚ ç›¸åŒçš„é‡åŒ–æ ¼å¼åœ¨ä¸åŒ GPU 上可能受益于不同的实现。Humming ä¼šä¸ºè®¾å¤‡é€‰æ‹©è§£ç ã€æ•°å€¼è½¬æ¢å’Œç¼©æ”¾è·¯å¾„ï¼Œå¹¶å°†å®ƒä»¬ä¸Žåˆé€‚çš„å† å­˜è®¿é—®å’Œè®¡ç®—æµæ°´çº¿é å¯¹ï¼Œä»¥å‡å°‘æŒ‡ä»¤å’Œæ•°æ®ç§»åŠ¨å¼€é”€ã€‚å ±äº«ç»„ä»¶ä½¿è¿™äº›ä¼˜åŒ–å¯åœ¨å ¼å®¹çš„ç¨ å¯†å’Œ MoE å† æ ¸ä¸­å¤ç”¨ã€‚

丰富的调优控制与设备特定启发式。 Humming 提供 tile 形状、流水线阶段数、warp 特化、TMA、MMA 指令、CTA 并发和 Stream-K 调度等控制项。设备特定的启发式将 GEMM å½¢çŠ¶ä¸Žèµ„æºçº¦æŸç›¸ç»“åˆï¼Œä»¥å¹³è¡¡å¹¶è¡Œæ€§ã€æ•°æ®å¤ç”¨å’Œæµæ°´çº¿å¼€é”€ï¼Œå¹¶éšç€è¾“å ¥å¤§å°å˜åŒ–é€‰æ‹©ä¸åŒçš„æ‰§è¡Œç­–ç•¥ã€‚å¼€å‘è€ å¯ä»¥ç›´æŽ¥ä½¿ç”¨é»˜è®¤é ç½®ï¼Œæˆ–é’ˆå¯¹ç‰¹å®šæ¨¡åž‹å’Œéƒ¨ç½²åœºæ™¯å¯¹å ¶è¿›è¡Œè°ƒæ•´ã€‚

结果

广泛的设备支持

Humming 的统一架构实现了广泛的 GPU 支持。对于 NVIDIA GPU，Humming å ¨é¢æ”¯æŒä»Ž Turing 起的每一代架构，为多代硬件带来统一的量化计算框架。

统一的 JIT 架构也让 Humming å¯ä»¥è½»æ¾æ‰©å±•åˆ°å ¶ä»–å ¼å®¹ CUDA 编程的 GPU。Humming å·²å ¨é¢æ”¯æŒå¹³å¤´å“¥é•‡å²³ 810Eï¼Œå°†å ¶å ±äº«å† æ ¸ä¸Žè°ƒä¼˜èƒ½åŠ›å¸¦åˆ°æ›´å¤šç¡¬ä»¶å¹³å°ã€‚

下图展示了 Humming 在各设备上支持的高性能激活类型：这些类型能够å åˆ†åˆ©ç”¨è®¾å¤‡çš„å³°å€¼è®¡ç®—èƒ½åŠ›ã€‚å›¾ä¸­çœç•¥çš„éƒ¨åˆ†ç±»åž‹ä»å¯é€šè¿‡å ¶ä»–æŒ‡ä»¤è·¯å¾„è¿è¡Œã€‚ä¾‹å¦‚ï¼ŒB300（SM103）上的 UMMA 指令不再支持 INT8 计算，但更早的 MMA 指令仍可执行 INT8 GEMMï¼Œåªæ˜¯æ€§èƒ½ä¼šå¤§å¹ é™ä½Žã€‚

量化支持

Humming å¯¹é‡åŒ–ç±»åž‹çš„æ”¯æŒä¸»è¦å–å†³äºŽæ¿€æ´»ç±»åž‹ä¸Žæƒé‡ç±»åž‹ä¹‹é—´çš„å ¼å®¹æ€§ã€‚å ¶ç»Ÿä¸€çš„é‡åŒ–è¡¨ç¤ºä¸Žå† æ ¸æž¶æž„å è®¸å¤šç§æ•´æ•°å’Œæµ®ç‚¹æƒé‡æ ¼å¼ä¸Žä¸åŒç²¾åº¦çš„æ¿€æ´»é åˆä½¿ç”¨ï¼Œä¸ºæ¨¡åž‹ç²¾åº¦å’Œéƒ¨ç½²æ€§èƒ½æä¾›çµæ´»é€‰æ‹©ã€‚

下图展示了 Humming 支持的权重与激活类型组合。行表示权重类型，列表示激活类型，勾选表示支持的组合。结合上文的设备支持矩阵，它勾勒出各硬件架构上可用的高性能量化计算路径。

Humming è¿˜æ”¯æŒä¸€ç§æŒ‡æ•°ä½ä¸ºé›¶çš„ç‰¹æ®Šæµ®ç‚¹æ ¼å¼ï¼Œå ¶å¯è¡¨ç¤ºçš„å€¼ä¸Žç›¸åŒä½å®½çš„å¯¹ç§°æ•´æ•°ä¸€è‡´ã€‚ä¾‹å¦‚ï¼ŒæŽ’é™¤ -128 和 float8e0m7 的 INT8 与 -127 到 127 çš„æ•´æ•°è¡¨ç¤ºç›¸åŒã€‚è¯¥æ ¼å¼æƒé‡çš„æ¿€æ´»ç±»åž‹å ¼å®¹æ€§ä¸Žç›¸åŒä½å®½çš„ INT 权重相同。Humming ä¼šè‡ªåŠ¨å°†ç¬¦åˆæ¡ä»¶çš„å¯¹ç§°æ•´æ•°æƒé‡è½¬æ¢ä¸ºè¯¥æ ¼å¼ï¼Œä»¥å®žçŽ°æ›´å¿«çš„åé‡åŒ–ã€‚

社区还发现了 Blackwell Tensor Core 的两项隐藏能力 ：原生支持 float4e0m3 和 float8e3m4 ï¼Œè¿™ä¸¤ç§ç‰¹æ®Šæµ®ç‚¹æ ¼å¼æœªåœ¨ PTX ç±»åž‹ç³»ç»Ÿä¸­å ¬å¼€ã€‚Humming 在所有 Blackwell æž¶æž„ä¸Šæ”¯æŒè¿™äº›ç‰¹æ€§ï¼ŒåŒ æ‹¬æ•°æ®ä¸­å¿ƒçš„ SM100/SM103 以及面向工作站和个人计算的 SM120/SM121 è®¾å¤‡ã€‚ä¸¤ç§æ ¼å¼å‡å·²é›†æˆåˆ°å ¶ç»Ÿä¸€çš„é‡åŒ–ä¸Žè®¡ç®—æ¡†æž¶ä¸­ï¼Œä¸ºä½Žç²¾åº¦æŽ¨ç†æä¾›æ›´å¤šåŽŸç”Ÿè®¡ç®—é€‰é¡¹ã€‚Rubin 不支持这两种隐藏的 Tensor Core æ ¼å¼ï¼Œå› æ­¤å¯¹è¿™äº›åŽŸç”Ÿè®¡ç®—è·¯å¾„çš„æ”¯æŒä¸ä¼šå»¶ä¼¸åˆ° Rubin。

性能

ä»¥ä¸‹å† æ ¸å¾®åŸºå‡†æµ‹è¯•ç”± vLLM 社区运行，测量 M 覆盖 1 到 8192 之间每个 2 的幂时的 GEMM æ€§èƒ½ã€‚å† æ ¸å®Œå ¨é¢„çƒ­ï¼Œå¹¶åœ¨å†· L2 ç¼“å­˜ä¸‹ç‹¬ç«‹è®¡æ—¶ï¼Œæ ·æœ¬ä¹‹é—´çŸ­æš‚åœé¡¿ã€‚MoE åŠ é€Ÿæ¯”ä½¿ç”¨ Gateup 和 Down GEMM å»¶è¿Ÿä¹‹å’Œã€‚æ¯å¼ å›¾éƒ½æ ‡æ˜Žäº†å ¶åŸºçº¿ã€é‡åŒ–æ ¼å¼ã€çŸ©é˜µç»´åº¦å’Œè½¯ä»¶ç‰ˆæœ¬ã€‚

#### WnA16

WnA16 是量化推理中最通用、应用最广泛的方法之一。它将低比特权重与 FP16/BF16 æ¿€æ´»ç›¸ç»“åˆï¼Œåœ¨å‡å°‘æƒé‡å­˜å‚¨å’Œå† å­˜æµé‡çš„åŒæ—¶ï¼Œåˆ©ç”¨å¹¿æ³›å¯ç”¨çš„ 16 ä½è®¡ç®—èƒ½åŠ›ã€‚è¿™ä½¿å ¶é€‚ç”¨äºŽè®¸å¤š GPU 世代和模型部署场景。W4A16 ç³»åˆ—å°¤å ¶å—åˆ°ç¤¾åŒºå ³æ³¨ï¼Œä¹Ÿæ˜¯ Humming 优化工作的重点。

对于密集 W4A16，Humming åœ¨å¤§å¤šæ•°æµ‹è¯•é ç½®ä¸­ä¸Ž vLLM Marlin æŒå¹³æˆ–ä¼˜äºŽåŽè€ ï¼Œå¹¶åœ¨è‹¥å¹² Hopper 和 Blackwell 工作负载上取得显著提升。下图以 cuBLAS ä½œä¸ºå ±åŒåŸºçº¿ã€‚åœ¨å°æ‰¹é‡ä¸‹ï¼ŒHumming 在多个设备上带来约 3–4 å€åŠ é€Ÿã€‚éšç€æ‰¹é‡å¢žå¤§ï¼Œå®ƒæŒç»­é«˜æ•ˆåˆ©ç”¨ç¡¬ä»¶ï¼Œåœ¨ H200 大批量工作负载上达到约 Marlin 的两倍性能，在部分 B300 åœºæ™¯ä¸‹è¶ è¿‡å ¶ä¸‰å€ã€‚

这些收益也延伸到 MoE。Humming 为 token 分布在各专家上、且每个专家的工作量随负载变化的工作负载提供高效的 Indexed 和 Grouped 路径。在下图中，H20 和 H200 上的 MXFP4 W4A16 在若干中大批量场景下达到约 Marlin 的两倍性能。在 B300 上，MXINT4 和 NVFP4 è·¯å¾„åœ¨å¤§æ‰¹é‡ä¸‹å–å¾—æ›´å¤§æ”¶ç›Šï¼Œå ¶ä¸­ Grouped 路径达到约 5 倍。SM120 和 SM121 上的结果也显示，Humming 在一系列工作负载中优于 Marlin 及相应的 FlashInfer 实现。

Humming 的优化还从 W4A16 æ‰©å±•åˆ°å ¶ä»– WnA16 ç»„åˆã€‚å ¶å ±äº«çš„åé‡åŒ–ä¸Žè®¡ç®—è·¯å¾„ä¸ºä¸åŒä½å®½çš„æ•´æ•°å’Œæµ®ç‚¹æƒé‡æä¾›å¼ºåŠ²æ€§èƒ½ï¼ŒåŒ æ‹¬ INT1、INT2、INT3 以及 FP4 到 FP8ã€‚ä¸‹å›¾å±•ç¤ºäº†è¿™äº›æ ¼å¼åœ¨å°æ‰¹é‡å’Œä¸­ç­‰æ‰¹é‡ä¸‹çš„å¹¿æ³›æ”¶ç›Šï¼Œå ¶ä¸­ SM120 ä¸Šçš„æœ€ä½Žä½å®½é ç½®è¶ è¿‡ BF16 cuBLAS 性能的 5 倍。在 H20 的大批量下，多个位宽保持接近 BF16 cuBLAS 的计算效率。这使模型能够在精度、存储和吞吐之间做出更细粒度的选择。

#### W8A8 / W4A8 / W4A4

在 Hopper 上，采用块级缩放的 FP8 W8A8 已成为高性能量化推理的主流方法。Humming 通过协同优化数据移动、缩放和专家调度，进一步改进这一成熟的计算路径。下方左列展示了 H20 和 H200 上的 W8A8 结果。在小批量和中等批量下，Humming 分别达到 DeepGEMM Contiguous 约 2.6 倍和 2.5 倍的性能，Indexed 和 Grouped ä¸¤ç§å¸ƒå±€å‡å ·æœ‰ç«žäº‰åŠ›ã€‚

随着更多模型采用原生 MXFP4 æƒé‡é‡åŒ–ï¼Œå°†æ›´å°çš„æƒé‡å ç”¨è½¬åŒ–ä¸ºå®žé™ æŽ¨ç†æ”¶ç›Šå˜å¾—è¶Šæ¥è¶Šæœ‰ä»·å€¼ã€‚Humming 通过高效的 W4A8 è·¯å¾„è§£å†³è¿™ä¸€é—®é¢˜ï¼šåœ¨åŠ è½½ MXFP4 权重时快速反量化，并使用 FP8 Tensor Core 进行计算，从而在 Hopper 上实现高效的 W4A8-MXFP4 执行。上方右列以 DeepGEMM W8A8 为基线。在小批量和中等批量下，Humming 的 W4A8 在 H20 上达到约 4.2 倍基线性能，在 H200 上达到 4.5 å€ï¼ŒåŒæ—¶ä¹Ÿå¤§å¹ ä¼˜äºŽåŒä¸€å›¾ä¸­ Humming 的 W8A8 结果。权重流量的减少转化为更低的 GEMM 延迟。FlashInfer CUTLASS W4A8 也是基于 Humming 方法的实现，而 Humming 当前的实现在此处展示的基准测试中取得更好性能。

Blackwell 为 MXFP8 和 MXFP4 ç­‰å¾®ç¼©æ”¾æ ¼å¼å¢žåŠ äº†åŽŸç”Ÿ Tensor Core 支持。从数据中心 SM100/SM103 设备到面向工作站和个人计算的 SM120/SM121 设备，Humming åˆ©ç”¨å„æž¶æž„çš„ä½Žç²¾åº¦è®¡ç®—èƒ½åŠ›ï¼Œåœ¨å•ä¸€å† æ ¸å’Œè°ƒä¼˜æ¡†æž¶å† æ”¯æŒ W8A8、W4A8 和 W4A4。

在 B300（SM103）上，Humming åœ¨å ¨éƒ¨ä¸‰ç§ MX æ ¼å¼ç»„åˆä¸‹å‡å±•çŽ°å‡ºæ˜¾è‘—çš„æ€§èƒ½æå‡ã€‚ä¸‹æ–¹æ¯ä¸ªé¢æ¿å‡ä»¥å¯¹åº”ç²¾åº¦ä¸‹çš„ DeepGEMM Contiguous 作为基线。在中小批量规模下，W8A8、W4A8 和 W4A4 分别达到约 1.25×、1.6× 和 1.4× 的基线性能，并在许多工作负载下优于对应的 FlashInfer 实现。Indexed 和 Grouped 两条路径均受益于原生低精度计算，为不同的专家布局提供了高效选项。

è¿™äº›æ”¶ç›ŠåŒæ ·å»¶ä¼¸è‡³ SM120 和 SM121。在下方结果中，Humming 的 MXFP4 × MXFP8 W4A8 和 NVFP4 × NVFP4 W4A4 在大多数测试点上优于 FlashInfer CUTLASS。在 SM121 的小批量规模下，W4A8 的性能提升约达 19%，W4A4 约达 9%；在 SM120 上，Humming åœ¨å¤§å¤šæ•°æµ‹è¯•è§„æ¨¡ä¸‹åŒæ ·é¢†å ˆã€‚ç»Ÿä¸€æž¶æž„å°†è¿™äº›ä½Žç²¾åº¦æŽ¨ç†ä¼˜åŒ–å¸¦åˆ°æ›´å¤šè®¾å¤‡ï¼ŒæŠŠé«˜æ•ˆé‡åŒ–è®¡ç®—æ‰©å±•åˆ°æ•°æ®ä¸­å¿ƒä¹‹å¤–ã€‚

#### 端到端推理性能

Humming 的 W4A8 æ”¶ç›Šå»¶ä¼¸è‡³æ¨¡åž‹æœåŠ¡ã€‚åœ¨æˆ‘ä»¬äºŽå «å— H20 GPU 上以 TP=8 进行的 DeepSeek-V4.1-Flash 测试中，保持模型 MXFP4 权重不变，并将 MoE 执行从 Marlin W4A16（BF16 激活）切换为 Humming W4A8（tokenwise FP8 激活），端到端吞吐提升约 14%–43% ，平均首 token 时延（TTFT）降低约 25%–32% 。吞吐优势随并发度增长而扩大：在并发度为 64 时，输出吞吐从 442.37 tokens/s 提升至 633.94 tokens/s，而平均每输出 token 耗时（TPOT）下降约 31% 。

并发度

MoE 后端

è¾“å ¥åžåï¼ˆtokens/s）↑

输出吞吐（tokens/s）↑

TTFT（ms）↓

TPOT（ms）↓

1

vLLM Marlin W4A16

2827.15

88.30

543.13

9.24

Humming W4A8

3226.23 (+14.1%)

100.77 (+14.1%)

383.80 (-29.3%)

8.46 (-8.5%)

2

4781.06

149.33

794.31

10.33

5646.98 (+18.1%)

176.38 (+18.1%)

555.66 (-30.0%)

9.20 (-10.9%)

4

7427.64

231.99

1422.32

11.72

9160.81 (+23.3%)

286.12 (+23.3%)

971.74 (-31.7%)

10.22 (-12.9%)

8

9190.03

287.04

2049.81

19.90

12263.89 (+33.4%)

383.04 (+33.4%)

1386.91 (-32.3%)

15.49 (-22.2%)

16

11124.67

348.96

2463.69

36.28

15171.00 (+36.4%)

475.89 (+36.4%)

1844.71 (-25.1%)

26.44 (-27.1%)

32

12669.43

398.83

3546.74

66.44

17804.79 (+40.5%)

560.49 (+40.5%)

2465.47 (-30.5%)

47.48 (-28.5%)

64

13937.42

442.37

5467.35

123.25

19973.00 (+43.3%)

633.94 (+43.3%)

3882.94 (-29.0%)

85.70 (-30.5%)

用法

安è£

使用 pip å®‰è£ Humming：

pip
install
humming-kernels

Humming 使用 CUDA ç¼–è¯‘å™¨æŒ‰éœ€ç”Ÿæˆå† æ ¸ã€‚å¦‚æžœç³»ç»Ÿä¸Šæ²¡æœ‰å¯ç”¨çš„ CUDA ç¼–è¯‘å™¨ï¼Œè¯·é€‰æ‹©ä»¥ä¸‹å®‰è£ é€‰é¡¹ä¹‹ä¸€ï¼Œä»¥åŒ å«é€‚ç”¨äºŽä½ æ‰€åœ¨çŽ¯å¢ƒ CUDA 主版本的编译依赖：

pip
install
"humming-kernels[cu12]"
# CUDA 12
pip
install
"humming-kernels[cu13]"
# CUDA 13

Humming 本身极为轻量，依赖简单，wheel åŒ ä» çº¦ 500 KBã€‚æŒ‰éœ€ç¼–è¯‘é¿å äº†æ‰“åŒ å¤§é‡é¢„ç¼–è¯‘å† æ ¸ã€‚åŽŸåˆ™ä¸Šï¼Œä»»ä½•èƒ½å¤Ÿè¿è¡Œ PyTorch 2.7 的环境也能运行 Hummingï¼Œå› æ­¤æ˜“äºŽé›†æˆåˆ°çŽ°æœ‰æŽ¨ç†çŽ¯å¢ƒä¸­ã€‚

在 vLLM 中使用 Humming

Humming 已深度集成到 vLLM 中。从 vLLM 0.29.0 开始，CUDA å®‰è£ å·²åŒ å« Humming ä½œä¸ºä¾èµ–ï¼Œå› æ­¤æ— éœ€å•ç‹¬å®‰è£ Humming。Humming 支持从 Turing 起所有架构的 NVIDIA GPU（SM75 及以上）。

--linear-backend 和 --moe-backend 选项分别选择 Humming ä½œä¸ºç¨ å¯† GEMM 和 MoE GEMM çš„è®¡ç®—åŽç«¯ã€‚ä¾‹å¦‚ï¼Œåœ¨å¼ é‡å¹¶è¡Œè§„æ¨¡ä¸º 8 æ—¶åŒæ—¶ä¸ºä¸¤è€ å¯ç”¨ Humming：

vllm
serve
deepseek-ai/DeepSeek-V4.1-Flash-1017
\
--tensor-parallel-size
8
\
--linear-backend
humming
--moe-backend
humming

--quantization humming é€‰é¡¹ä» ä½œä¸ºå ¼å®¹è·¯å¾„ï¼Œç”¨äºŽ Humming 支持但 vLLM ä¸å†æ”¯æŒæˆ–å°šæœªåŽŸç”Ÿæ”¯æŒçš„æ ¼å¼ï¼Œä¾‹å¦‚ GPTQ INT3ã€‚ä¾‹å¦‚ï¼Œé€šè¿‡æ­¤å ¼å®¹è·¯å¾„åŠ è½½ GPTQ INT3 检查点：

vllm
serve
hfl/Qwen2.5-VL-3B-Instruct-GPTQ-Int3
--quantization
humming

选择激活类型

é»˜è®¤æƒ å†µä¸‹ï¼ŒHumming 使用 checkpoint ä¸­æŒ‡å®šçš„æ¿€æ´»ç±»åž‹ï¼Œå¹¶æ ¹æ®è®¾å¤‡è°ƒæ•´æ‰§è¡Œæ–¹å¼ã€‚å¦‚æžœ GPU 不支持该计算类型，Humming ä¼šå›žé€€åˆ°å—æ”¯æŒçš„æ›´é«˜ç²¾åº¦ç±»åž‹ã€‚ä¾‹å¦‚ï¼Œä¸å ·å¤‡åŽŸç”Ÿ MXFP4 支持的设备可以自动使用 FP8、BF16 或 FP16 计算路径，从而让同一个量化 checkpoint è·¨ GPU ä»£é™ è¿è¡Œã€‚

用户可以通过 vLLM çš„é‡åŒ–é ç½®æŽ¥å£è¦†ç›–é»˜è®¤è®¾ç½®ï¼Œä¸ºéƒ¨ç½²é€‰æ‹©æ¿€æ´»ç±»åž‹ã€‚ä»¥ä¸‹ç¤ºä¾‹å°† MoE 层的激活类型设置为 MXFP8：

vllm
serve
openai/gpt-oss-20b
--moe-backend
humming
--quantization-config.moe.activation
mxfp8

未来计划

æ›´å¹¿æ³›ã€æ›´æ·±å ¥çš„è®¾å¤‡æ”¯æŒ

Humming å°†ç»§ç»­æ‰©å¤§ç¡¬ä»¶è¦†ç›–èŒƒå›´ï¼ŒåŒæ—¶é€‚é æ–°è®¾å¤‡çš„è®¡ç®—ä¸Žå† å­˜èƒ½åŠ›ï¼Œå¸®åŠ©æ›´å¤šå¹³å°é‡Šæ”¾é‡åŒ–æŽ¨ç†çš„æ€§èƒ½æ½œåŠ›ã€‚ä¼˜å ˆäº‹é¡¹åŒ æ‹¬ï¼š

• 支持更多 Rubin æž¶æž„ç‰¹æ€§ï¼ŒåŒ æ‹¬ Locality Domainsï¼Œä»¥è¿›ä¸€æ­¥ä¼˜åŒ–ä»»åŠ¡æ˜ å°„å’Œæ•°æ®è®¿é—®ã€‚

• 将支持扩展到更多 T-Head PPU è®¾å¤‡ï¼ŒåŒ æ‹¬ Zhenwu M890。

• æ”¯æŒæ›´å¤šå ¼å®¹ CUDA ç¼–ç¨‹çš„è®¾å¤‡ï¼Œå°†çŽ°æœ‰é‡åŒ–æ ¼å¼ã€å† æ ¸ç»„ä»¶å’Œè°ƒä¼˜èƒ½åŠ›å¸¦åˆ°æ›´å¤šç¡¬ä»¶å¹³å°ã€‚

åŠ å ¥ Autotune

通过针对性调优，Humming åœ¨å¤§å¤šæ•°åœºæ™¯ä¸‹å·²ç»èƒ½å¤Ÿè¾¾åˆ°ä¸šç•Œé¢†å ˆçš„æ€§èƒ½ã€‚ç„¶è€Œï¼Œæ¨¡åž‹å½¢çŠ¶ã€é‡åŒ–é ç½®å’Œè®¾å¤‡èµ„æºçš„èŒƒå›´ååˆ†å¹¿æ³›ï¼Œå¯å‘å¼æ–¹æ³•éš¾ä»¥åœ¨æ¯ç§æƒ å†µä¸‹éƒ½é€‰å‡ºæœ€ä¼˜é ç½®ã€‚

å› æ­¤ï¼ŒHumming è®¡åˆ’å¼•å ¥ autotuneï¼Œåœ¨ç›®æ ‡è®¾å¤‡ä¸Šå¯¹å€™é€‰é ç½®è¿›è¡ŒåŸºå‡†æµ‹è¯•ï¼Œä¸ºæ¯ä¸ªå·¥ä½œè´Ÿè½½æ‰¾åˆ°æ›´å¥½çš„æ‰§è¡Œå‚æ•°ã€‚Autotune 将补å çŽ°æœ‰å¯å‘å¼æ–¹æ³•ï¼Œå‡å°‘æ‰‹åŠ¨é ç½®æœç´¢ï¼Œå¸®åŠ©æ›´å¤šå·¥ä½œè´Ÿè½½æŽ¥è¿‘å† æ ¸å¯è¾¾åˆ°çš„æœ€ä½³æ€§èƒ½ã€‚

与推理框架的更深集成

Humming 将继续深化与 vLLM ç­‰æŽ¨ç†æ¡†æž¶çš„é›†æˆï¼Œæ‰©å¤§åœ¨é‡åŒ–æ¨¡åž‹åŠ è½½ã€åŽç«¯é€‰æ‹©ã€å¹¶è¡ŒæŽ¨ç†ã€è¿è¡Œæ—¶è°ƒåº¦å’Œè®¾å¤‡é€‚é ç­‰æ–¹é¢çš„æ”¯æŒã€‚æˆ‘ä»¬çš„ç›®æ ‡æ˜¯è®©æ›´å¤šé‡åŒ–æ–¹æ³•èƒ½å¤Ÿé åˆæ¡†æž¶çš„å®Œæ•´èƒ½åŠ›ï¼ŒåŒ æ‹¬é¢å‘ä¸“å®¶å¹¶è¡Œï¼ˆEP）的量化 All-to-All（A2Aï¼‰é€šä¿¡ç­‰ä¼˜åŒ–ã€‚è¿™äº›å·¥ä½œå°†ä½¿é«˜æ€§èƒ½é‡åŒ–å† æ ¸åœ¨å¤šæ ·åŒ–éƒ¨ç½²åœºæ™¯ä¸­æ›´æ˜“ä½¿ç”¨ï¼Œå¹¶å¸®åŠ©å°†å† æ ¸æ€§èƒ½æ”¶ç›Šè½¬åŒ–ä¸ºç«¯åˆ°ç«¯æŽ¨ç†æ”¶ç›Šã€‚

致谢

我们感谢蚂蚁集团的 Jinzhen Lin、Zekai Gu 和 Zhilin Huang，以及 Michael Goin、Roberto L. Castro、Chauncey Jiang 和 vLLM ç¤¾åŒºçš„å ¶ä»–æˆå‘˜ï¼Œæ„Ÿè°¢ä»–ä»¬åœ¨ Humming 开发过程中的贡献与支持。

我们还要感谢 vLLM 社区提供开发、测试和 CI 资源，使我们能够在广泛的硬件和工作负载上验证并改进 Humming。

Humming 的设计与优化灵感来自 Marlin、vLLM Marlin、LMDeploy、DeepGEMM、FlashInfer、CUTLASS åŠå ¶ä»–å¼€æºé¡¹ç›®ã€‚æˆ‘ä»¬æ„Ÿè°¢è¿™äº›é¡¹ç›®çš„ä½œè€ å’Œè´¡çŒ®è€ åˆ†äº«ä»–ä»¬çš„æƒ³æ³•ä¸Žå®žçŽ°ï¼Œå¸®åŠ©æŽ¨åŠ¨æ•´ä¸ªç¤¾åŒºçš„é«˜æ€§èƒ½é‡åŒ–æŽ¨ç†å‘å±•ã€‚

来源： vLLM 官方博客 · vllm.ai

开源生态

部署工程

产品更新

#

开源/仓库

部署/工程