Humming wird angekündigt: Schnelle, flexible Kernels für quantisierte Inferenz

Ant Group und die vLLM-Community veröffentlichen Humming als Open Source, eine universelle, leistungsstarke Kernel-Bibliothek für quantisierte Inferenz. Humming bietet die folgenden Funktionen:

• Breite Geräteunterstützung: Unterstützt NVIDIA-GPU-Architekturen ab Turing sowie CUDA-kompatible Geräte wie T-Heads Zhenwu 810E.
• Umfassende Quantisierungsunterstützung: Eine einheitliche Architektur unterstützt vielfältige Kombinationen von Gewichts- und Aktivierungsformaten für dense- und MoE-Workloads und deckt eine breite Palette von Ganzzahl- und Gleitkommatypen, Bitbreiten und Skalierungsschemata ab.
• Starke Ausführungsleistung: Übertrifft bestehende Implementierungen von vLLM Marlin, FlashInfer und DeepGEMM bei vielen repräsentativen Workloads deutlich.

Repository: github.com/vllm-project/humming

Hintergrund

Mit wachsenden Modellen und Inferenz-Workloads werden Speicherkapazität, Speicherbandbreite und Rechenkosten zunehmend zum Engpass bei der Bereitstellung. Dies zeigt sich besonders bei MoE-Modellen: Jeder Token aktiviert nur eine Teilmenge der Experten, dennoch muss der vollständige Satz an Gewichten gespeichert werden. DeepSeek-V3 hat beispielsweise 671 Milliarden Parameter, von denen 37 Milliarden pro Token aktiviert werden. Quantisierung verringert die Präzision von Gewichten und Aktivierungen, um Speichernutzung und Datenbewegung zu reduzieren, und kann zudem die Recheneffizienz verbessern.

Die Anforderungen an die Quantisierung variieren je nach Modell und GPU-Generation. Übliche Konfigurationen sind WNA16, W8A8, W4A8 und W4A4, wobei W und A die Bitbreiten von Gewicht und Aktivierung angeben und N eine variable Gewichtsbitbreite bezeichnet. Selbst bei gleicher Präzision unterscheiden sich Schemata wie INT4 und NVFP4 in numerischem Format und Skalierung. Skalierungsgranularität, Skalendarstellung, Gruppengröße sowie symmetrische oder asymmetrische Quantisierung erweitern den Konfigurationsraum zusätzlich.

Diese Einsparungen bei Speicherung und Berechnung führen nur mit effizienten Kernels zu geringerer Latenz und höherem Durchsatz. Prefill, Decode und unterschiedliche MoE-Expertenauslastungen stellen unterschiedliche Anforderungen an Berechnung und Speicher. Kernels müssen sich an diese Workloads anpassen und dabei den Overhead für Konvertierung, Dequantisierung und Scheduling gering halten.

Die Entwicklung und Integration eines dedizierten Kernels für jede Kombination schafft einen wachsenden Wartungsaufwand. Jede Ergänzung kann eigene Gewichtsvorverarbeitung, Kernel-Auswahl, Framework-Adapter und Tests erfordern. Optimierungen müssen anschließend über mehrere Implementierungen hinweg gepflegt werden.

Humming löst dieses Problem, indem es Quantisierungsbeschreibungen von Ausführungswahl trennt und Kernels aus gemeinsamen Komponenten zusammensetzt.

Universelle Architektur

Hummings Ausführungspfad besteht aus drei Teilen: die quantisierte Berechnung darstellen, Konfigurationen für Hardware und Workload auswählen und einen Kernel aus CUDA-Komponenten kompilieren.

Eine einheitliche Quantisierungsdarstellung

Von verschiedenen Quantisierungswerkzeugen exportierte Modelle unterscheiden sich in numerischem Format, Parameterorganisation und Gewichtslayout. Humming konvertiert unterstützte Exportformate in das Humming-Standardformat: eine gemeinsame Quantisierungsbeschreibung mit konsistent organisierten Gewichten, Skalen, Nullpunkten und anderen Tensoren.

Diese Darstellung bewahrt die Quantisierungssemantik jedes Schemas und beschreibt Gewichtsformat, Skalierungsgranularität, Skalendatentyp, Nullpunkte und sekundäre Skalen als separate Parameter. Aktivierungen haben ihr eigenes numerisches Format und ihre eigenen Quantisierungseinstellungen. Diese Parameter können unter Einhaltung von Format- und Hardwareeinschränkungen kombiniert werden.

So können beispielsweise gruppweise INT4-Gewichte mit BF16-Aktivierungen gepaart werden. Gleitkommagewichte mit geringer Präzision und unterschiedlichen Skalierungsschemata verwenden dasselbe Parametersystem. Dies bietet Inferenz-Frameworks eine gemeinsame Schnittstelle für Modelle aus verschiedenen Quellen.

Tuning und JIT-Kompilierung für jeden Workload

Das Standardformat bietet eine gemeinsame Darstellung über Exportformate hinweg. Die Vorverarbeitung validiert anschließend Parameter, polstert Dimensionen auf und packt diese Tensoren in Layouts um, die für den Zielkernel geeignet sind. Humming kombiniert die Quantisierungsbeschreibung mit der Gewichtungsform, dem GEMM-Typ, der Ziel-GPU und den Berechnungsoptionen, um die Ausführung zu konfigurieren.

Drei Konfigurationen beschreiben den Kernel:

• Layer Config beschreibt Tensorformen, Datentypen, Quantisierungsparameter und die Zielarchitektur.
• Compute Config gibt den GEMM-Typ und Berechnungsoptionen an, wie Batch-Invarianz und das Layout der Eingabeskala.
• Tuning Config wählt Tile-Formen, Pipelinestufen-Anzahlen, Warp-Spezialisierung, Speicherzugriff und MMA-Instruktionen aus. Er steuert auch Nebenläufigkeit und Scheduling; TMA kann pro Tensor konfiguriert werden.

Standardmäßig wählt Humming Tuning-Konfigurationen anhand gerätespezifischer Heuristiken aus. Es kann unterschiedliche Konfigurationen für Bereiche von Eingabegrößen vorbereiten und zur Laufzeit an den passenden Kernel dispatchen.

Humming’s vereinheitlichte Architektur macht es unkompliziert, Ausführungsparameter zu tunen und Pipelines für verschiedene Workloads auszuwählen. Bei kleinen Batch-Größen kann eine mehrstufige Pipeline mit asynchronen Kopien (cp.async) und MMA eine Pipeline, die auf TMA und UMMA basiert, übertreffen. Bei größeren Batch-Größen können TMA und UMMA einen besseren Durchsatz bieten. Humming unterstützt beide Ausführungspfade innerhalb desselben Konfigurations- und Komponenten-Frameworks, sodass es für jeden Workload die effektivere Strategie auswählen kann.

Zusammen spezialisieren die drei Konfigurationen Kernel, die mit NVRTC oder NVCC kompiliert wurden. Kompilierte Kernel werden zur Wiederverwendung zwischengespeichert. Die Inferenz kombiniert den ausgewählten Kernel mit den vorverarbeiteten Gewichten und Laufzeit-Aktivierungen.

Gemeinsam genutzte CUDA-Komponenten über Formate und Architekturen hinweg

Auf CUDA-Ebene organisiert Humming Aufgabenplanung, Datenbewegung, Dequantisierung und Skalierung, Tensor-Core-Berechnung sowie Ausgabeverarbeitung in wiederverwendbaren Komponenten. Es bietet standardmäßige, warp-spezialisierte und dedizierte UMMA-Pipelines. Diese Pipelines teilen sich kompatible Komponenten und nutzen gleichzeitig den von jeder Ausführungsstrategie benötigten Datenfluss und die entsprechende Synchronisation.

Jede Gruppe von Komponenten behandelt eine andere Variationsquelle:

• Die numerische Verarbeitung behandelt Quantisierungsformate. Dekodier- und Arithmetikkomponenten berücksichtigen Bitbreite, Ganzzahl- oder Gleitkommadarstellung, Skalen und Nullpunkte. Dequantisierung und Skalierung sind in die Berechnung des Kernels integriert.
• Speicher- und Compute-Backends behandeln GPU-Architekturen. Humming wählt Datenbewegungs- und Tensor-Core-Implementierungen für die Zielarchitektur aus und nutzt asynchrone Kopien, TMA und Multicast, wo unterstützt.
• Task-Zuordnung und -Planung behandeln dichte und MoE-Workloads. Jeder GEMM-Typ verwendet die geeignete Planungslogik. Skalierung und Bias-Addition können ebenfalls in die Ausgabestufe integriert werden.

Die JIT-Kompilierung kombiniert diese Komponenten gemäß den drei Konfigurationen und entfernt ungenutzte Implementierungszweige. Ein neues Format oder eine neue Architektur kann daher durch Erweiterung der relevanten Komponenten unterstützt werden. Verbesserungen an einer gemeinsam genutzten Komponente kommen kompatiblen Konfigurationen zugute und verringern den doppelten Entwicklungs- und Wartungsaufwand separater spezialisierter Kernel.

Diese CUDA-Implementierung ist leichtgewichtig und vermeidet schwergewichtige Kernel-Framework-Abhängigkeiten, was auch die JIT-Kompilierung schnell hält. In unserem einzelnen W4A16-Kernel-Kompilierungstest auf einem H200-Server mit CUDA 13.4 sank die Kompilierungszeit von 23,35 Sekunden für CUTLASS mit NVCC auf 1,82 Sekunden für Humming mit NVCC oder 1,88 Sekunden mit NVRTC, wobei die Kompilierungscaches deaktiviert waren. Dies reduziert die Kosten der Kernel-Vorbereitung für neue Workloads und macht iteratives Tuning praktikabler. Humming kompiliert Kernel während des Modell-Warmups auch parallel. In unseren Tests dauert das Kompilieren und Laden der Kernel, die alle unterstützten M-Größen für den GEMM einer einzelnen Schicht abdecken, nur etwa 2–4 Sekunden, sodass es praktikabel ist, mehrere spezialisierte Kernel vor der Inferenz vorzubereiten.

Leistungsoptimierung basierend auf Geräteparametern

Humming optimiert Kernel tiefgehend unter Verwendung sowohl der GEMM-Konfiguration als auch von Geräteparametern. Quantisierungsformate, Matrixformen, GPU-Instruktionsunterstützung, SM-Anzahl, Shared Memory und Registerressourcen fließen alle in die Ausführungswahl ein und verwandeln die Vorteile der Low-Bit-Speicherung in praktische Inferenzleistung.

Effiziente W4A8-MXFP4-Inferenz auf bestehenden GPUs. Architekturen vor Blackwell verfügen über keine native MXFP4-Tensor-Core-Unterstützung, und konventionelle W4A8-MXFP4-Kernel können bei größeren M durch den Overhead von Dequantisierung und Skalierung begrenzt sein. Humming löst dies mit einem schnellen MXFP4-zu-INT8/FP8-Dequantisierungspfad, der MXFP4-Gewichte innerhalb des Kernels in ein unterstütztes Berechnungsformat umwandelt. In diesem Setting übertrifft der Pfad konventionelle W4A8-Implementierungen bei kleinem M erheblich und erreicht bei großem M nahezu W8A8-Leistung. GPUs wie H200 können daher hochperformante W4A8-MXFP4-Kernel ausführen und gleichzeitig die Speichervorteile von 4-Bit-Gewichten beibehalten.

An jedes Gerät angepasste Dequantisierungspfade. Dasselbe Quantisierungsformat kann auf verschiedenen GPUs von unterschiedlichen Implementierungen profitieren. Humming wählt Dekodierungs-, Zahlenkonvertierungs- und Skalierungspfade für das Gerät aus und kombiniert sie mit geeigneten Speicherzugriffs- und Compute-Pipelines, um Instruktions- und Datenbewegungsaufwand zu reduzieren. Gemeinsam genutzte Komponenten machen diese Optimierungen über kompatible dichte und MoE-Kernel hinweg wiederverwendbar.

Umfangreiche Tuning-Steuerung mit gerätespezifischen Heuristiken. Humming legt Tile-Formen, Pipeline-Stufenanzahlen, Warp-Spezialisierung, TMA, MMA-Instruktionen, CTA-Konkurrenz und Stream-K-Planung offen. Gerätespezifische Heuristiken kombinieren GEMM-Formen mit Ressourcenbeschränkungen, um Parallelismus, Datenwiederverwendung und Pipeline-Overhead auszubalancieren und bei sich ändernden Eingabegrößen unterschiedliche Ausführungsstrategien zu wählen. Entwickler können die Standardkonfigurationen direkt verwenden oder sie für bestimmte Modelle und Bereitstellungsszenarien verfeinern.

Ergebnisse

Breite Geräteunterstützung

Hummings einheitliche Architektur ermöglicht breite GPU-Unterstützung. Für NVIDIA-GPUs unterstützt Humming vollständig jede Architektur ab Turing und bringt ein gemeinsames quantisiertes Rechenframework über mehrere Hardwaregenerationen.

Die vereinheitlichte JIT-Architektur macht es auch einfach, Humming auf andere mit CUDA-Programmierung kompatible GPUs zu erweitern. Humming unterstützt T-Heads Zhenwu 810E bereits vollständig und bringt seine gemeinsamen Kernel und Tuning-Fähigkeiten auf weitere Hardwareplattformen.

Die folgende Abbildung zeigt die von Humming auf jedem Gerät unterstützten hochperformanten Aktivierungstypen: Typen, die die Spitzenrechenleistung des Geräts voll ausnutzen können. Einige in der Abbildung nicht aufgeführte Typen können weiterhin über andere Instruktionspfade ausgeführt werden. Beispielsweise unterstützen UMMA-Instruktionen auf dem B300 (SM103) keine INT8-Berechnung mehr, aber frühere MMA-Instruktionen können weiterhin INT8-GEMM ausführen, mit deutlich geringerer Leistung.

Quantisierungsunterstützung

Die Unterstützung von Humming für Quantisierungstypen hängt in erster Linie von der Kompatibilität zwischen Aktivierungs- und Gewichtstypen ab. Die vereinheitlichte Quantisierungsdarstellung und Kernelarchitektur ermöglichen es, dass eine breite Palette von Ganzzahl- und Gleitkomma-Gewichtsformaten mit Aktivierungen unterschiedlicher Präzision zusammenarbeiten, was flexible Auswahlmöglichkeiten für Modellpräzision und Deployment-Leistung bietet.

Die folgende Abbildung zeigt die von Humming unterstützten Kombinationen aus Gewichts- und Aktivierungstypen. Die Zeilen repräsentieren Gewichtstypen, die Spalten Aktivierungstypen, und Häkchen kennzeichnen unterstützte Kombinationen. Zusammen mit der obenstehenden Geräteunterstützungsmatrix skizziert sie die auf jeder Hardwarearchitektur verfügbaren leistungsstarken quantisierten Berechnungspfade.

Humming unterstützt außerdem ein spezielles Gleitkommaformat mit null Exponentenbits, dessen darstellbare Werte denen symmetrischer Ganzzahlen bei derselben Bitbreite entsprechen. Beispielsweise stellen INT8 ohne -128 und float8e0m7 beide Ganzzahlen von -127 bis 127dar. Gewichte in diesem Format haben dieselbe Kompatibilität mit Aktivierungstypen wie INT-Gewichte derselben Bitbreite. Humming konvertiert geeignete symmetrische Ganzzahl-Gewichte automatisch in dieses Format für eine schnellere Dequantisierung.

Die Community hat außerdem zwei versteckte Fähigkeiten der Blackwell Tensor Cores aufgedeckt: native Unterstützung für float4e0m3 und float8e3m4, zwei spezielle Gleitkommaformate, die im PTX-Typsystem nicht offengelegt werden. Humming unterstützt diese Funktionen auf allen Blackwell-Architekturen, einschließlich der Datacenter-Geräte SM100/SM103 und der Geräte SM120/SM121 für Workstations und Personal Computing. Beide Formate sind in sein einheitliches Quantisierungs- und Rechenframework integriert und bieten damit mehr native Rech Optionen für Inferenz mit geringer Präzision. Rubin unterstützt diese beiden versteckten Tensor-Core-Formate nicht, sodass die Unterstützung für diese nativen Rechenpfade nicht auf Rubin ausgedehnt wird.

Leistung

Die folgenden Kernel-Mikrobenchmarks, durchgeführt von der vLLM-Community, messen die GEMM-Leistung mit M, das jede Zweierpotenz von 1 bis 8192 abdeckt. Die Kernels sind vollständig aufgewärmt und werden unabhängig voneinander mit einem kalten L2-Cache gemessen, mit einer kurzen Pause zwischen den Messungen. Die MoE-Beschleunigungen verwenden die Summe der Gateup- und Down-GEMM-Latenzen. Jede Abbildung nennt ihre Baseline, Quantisierungsformate, Matrixdimensionen und Softwareversionen.

WnA16

WnA16 ist einer der vielseitigsten und am weitesten verbreiteten Ansätze für quantisierte Inferenz. Er kombiniert Gewichte mit niedriger Bitbreite mit FP16/BF16-Aktivierungen, wodurch der Gewichte-Speicherbedarf und der Speicherverkehr reduziert werden, während weit verfügbare 16-Bit-Rechenfähigkeiten genutzt werden. Dies macht ihn für viele GPU-Generationen und Modellbereitstellungsszenarien geeignet. Die W4A16-Familie hat besonderes Interesse in der Community geweckt und ist ein zentraler Fokus der Optimierungsarbeit an Humming.

Bei dense W4A16 entspricht Humming in den meisten getesteten Konfigurationen dem Niveau von vLLM Marlin oder übertrifft es, mit deutlichen Gewinnen bei mehreren Hopper- und Blackwell-Workloads. Die folgende Abbildung verwendet cuBLAS als gemeinsame Baseline. Bei kleinen Batch-Größen liefert Humming auf mehreren Geräten etwa 3–4× Beschleunigungen. Mit wachsender Batch-Größe nutzt es die Hardware weiterhin effizient und erreicht etwa die doppelte Leistung von Marlin bei Large-Batch-H200-Workloads und mehr als das Dreifache in einigen B300-Fällen.

Diese Gewinne erstrecken sich auch auf MoE. Humming bietet effiziente Indexed- und Grouped-Pfade für Workloads, bei denen Tokens auf Experten verteilt sind und der Arbeitsumfang pro Expert mit der Last variiert. In der folgenden Abbildung erreicht MXFP4 W4A16 auf H20 und H200 in mehreren mittel- bis großformatigen Batch-Fällen etwa die doppelte Leistung von Marlin. Auf B300 erzielen die MXINT4- und NVFP4-Pfade bei großen Batch-Größen noch größere Gewinne, wobei der Grouped-Pfad etwa 5× erreicht. Ergebnisse auf SM120 und SM121 zeigen ebenfalls, dass Humming Marlin und die entsprechenden FlashInfer-Implementierungen über eine Reihe von Workloads hinweg übertrifft.

Humming's Optimierungen erstrecken sich auch über W4A16 hinaus auf andere WnA16-Kombinationen. Seine gemeinsamen Dequantisierungs- und Rechenpfade liefern starke Leistung für Ganzzahl- und Gleitkommagewichte mit unterschiedlichen Bitbreiten, einschließlich INT1, INT2, INT3 und FP4 bis FP8. Die folgende Abbildung zeigt breite Gewinne über diese Formate bei kleinen und mittleren Batch-Größen hinweg, wobei die Konfigurationen mit der niedrigsten Bitbreite auf SM120 mehr als 5× die Leistung von BF16 cuBLAS überschreiten. Bei großen Batch-Größen auf H20 halten mehrere Bitbreiten die Recheneffizienz nahe an BF16 cuBLAS. Dies gibt Modellen feinere Auswahlmöglichkeiten zwischen Präzision, Speicher und Durchsatz.

W8A8 / W4A8 / W4A4

Mit Hopper wurde FP8 W8A8 mit blockweisem Scaling zu einem etablierten Ansatz für leistungsstarke quantisierte Inferenz. Humming verbessert diesen bewährten Rechenpfad weiter, indem Datenbewegung, Scaling und Experten-Planung gemeinsam optimiert werden. Die linke Spalte unten zeigt W8A8-Ergebnisse auf H20 und H200. Bei kleinen und mittleren Batch-Größen erreicht Humming etwa 2.6× bzw. 2.5× die Leistung von DeepGEMM Contiguous, mit konkurrenzfähiger Leistung sowohl bei Indexed- als auch bei Grouped-Layouts.

Da immer mehr Modelle native MXFP4-Gewichtsquantisierung übernehmen, wird es zunehmend wertvoll, den kleineren Gewichtsfußabdruck in praktische Inferenzgewinne umzuwandeln. Humming löst dies mit einem effizienten W4A8-Pfad, der MXFP4-Gewichte beim Laden schnell dequantisiert und FP8 Tensor Cores für die Berechnung verwendet, was eine effiziente W4A8-MXFP4-Ausführung auf Hopper ermöglicht. Die rechte Spalte oben verwendet DeepGEMM W8A8 als Baseline. Bei kleinen und mittleren Batch-Größen erreicht Hummings W4A8 etwa 4.2× der Baseline-Leistung auf H20 und 4.5× auf H200 und übertrifft damit auch deutlich Hummings W8A8-Ergebnisse in derselben Abbildung. Reduzierter Gewichtsverkehr übersetzt sich in niedrigere GEMM-Latenz. FlashInfer CUTLASS W4A8 ist ebenfalls eine Implementierung auf Basis von Hummings Ansatz, während Hummings aktuelle Implementierung in den hier gezeigten Benchmarks eine bessere Leistung erzielt.

Blackwell fügt native Tensor-Core-Unterstützung für Microscaling-Formate wie MXFP8 und MXFP4 hinzu. Von den Datacenter-Geräten SM100/SM103 bis zu den Geräten SM120/SM121 für Workstations und Personal Computing nutzt Humming die Rechenfähigkeiten mit geringer Präzision jeder Architektur, um W8A8, W4A8 und W4A4 innerhalb eines Kernels und Tuning-Frameworks zu unterstützen.

Auf B300 (SM103) zeigt Humming erhebliche Leistungsgewinne über alle drei MX-Formatkombinationen. Jedes der folgenden Diagramme verwendet DeepGEMM Contiguous bei der jeweiligen Präzision als Basislinie. Bei kleinen und mittleren Batch-Größen erreichen W8A8, W4A8 und W4A4 jeweils approximately 1.25×, 1.6× bzw. 1.4× der Basislinienleistung und übertreffen die entsprechenden FlashInfer-Implementierungen bei vielen Workloads. Sowohl die Indexed- als auch die Grouped-Pfade profitieren von nativer Berechnung in niedriger Präzision und bieten effiziente Optionen für verschiedene Expert-Layouts.

Diese Gewinne erstrecken sich auch auf SM120 und SM121. In den unten gezeigten Ergebnissen übertreffen Hummings MXFP4 × MXFP8 W4A8 und NVFP4 × NVFP4 W4A4 FlashInfer CUTLASS an den meisten getesteten Punkten. Auf SM121 erreichen die Leistungsgewinne bei kleinen Batch-Größen etwa 19% für W4A8 und 9% für W4A4; auf SM120 liegt Humming ebenfalls bei den meisten getesteten Größen vorn. Die einheitliche Architektur bringt diese Optimierungen für Inferenz mit geringer Präzision auf mehr Geräte und erweitert effiziente quantisierte Berechnung über das Rechenzentrum hinaus.

End-to-End-Inferenzleistung

Humming W4A8 bringt seine Vorteile auch beim Modell-Serving. In unserem DeepSeek-V4.1-Flash-Test auf acht H20-GPUs mit TP=8 verbessert sich der End-to-End-Durchsatz um ungefähr, wenn die MXFP4-Gewichte des Modells unverändert bleiben und die MoE-Ausführung von Marlin W4A16 (BF16-Aktivierungen) auf Humming W4A8 (tokenweise FP8-Aktivierungen) umgestellt wird. 14%–43% und reduziert die mittlere Zeit bis zum ersten Token (TTFT) um ungefähr 25%–32%. Der Durchsatzvorteil wächst mit der Konkurrenzlast: Bei einer Konkurrenz von 64 steigt der Output-Durchsatz von 442,37 auf 633,94 tokens/s, während die mittlere Zeit pro Output-Token (TPOT) um approximately sinkt. 31%.

Verwendung

Installation

Installiere Humming mit pip:

pip install humming-kernels

Humming verwendet einen CUDA-Compiler, um Kernels bei Bedarf zu generieren. Falls auf dem System kein CUDA-Compiler verfügbar ist, wählen Sie eine der folgenden Installationsoptionen, um die Kompilierungsabhängigkeiten für die CUDA-Hauptversion Ihrer Umgebung einzubinden:

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

Humming selbst ist extrem leichtgewichtig, mit einfachen Abhängigkeiten und einem Wheel-Paket von nur etwa 500 KB. Die Kompilierung bei Bedarf vermeidet das Mitliefern einer großen Sammlung vorkompilierter Kernel. Im Prinzip kann jede Umgebung, die PyTorch 2.7 ausführen kann, auch Humming ausführen, was die Integration in bestehende Inferenzumgebungen erleichtert.

Humming in vLLM verwenden

Humming ist tief in vLLM integriert. Ab vLLM 0.29.0 enthalten CUDA-Installationen Humming als Abhängigkeit, sodass keine separate Humming-Installation mehr nötig ist. Humming unterstützt NVIDIA-GPUs aller Architekturen ab Turing (SM75 und höher).

Die --linear-backend und --moe-backend Optionen wählen Humming als Compute-Backend für dense GEMM bzw. MoE GEMM. Zum Beispiel aktiviert Folgendes Humming für beide mit einer tensor-parallel size von 8:

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

Die --quantization humming Option ist ausschließlich als Kompatibilitätspfad für Formate gedacht, die Humming unterstützt, vLLM aber nicht mehr oder noch nicht nativ unterstützt, wie etwa GPTQ INT3. Beispielsweise lädt man einen GPTQ INT3-Checkpoint über diesen Kompatibilitätspfad:

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

Auswahl von Aktivierungstypen

Standardmäßig verwendet Humming den im Checkpoint angegebenen Aktivierungstyp und passt die Ausführung an das Gerät an. Wenn die GPU diesen Rechentyp nicht unterstützt, fällt Humming auf einen unterstützten Typ mit höherer Präzision zurück. Beispielsweise können Geräte ohne native MXFP4-Unterstützung automatisch FP8-, BF16- oder FP16-Rechenpfade verwenden, sodass derselbe quantisierte Checkpoint über GPU-Generationen hinweg laufen kann.

Benutzer können den Standard über vLLMs Quantisierungskonfigurationsschnittstelle außer Kraft setzen, um einen Aktivierungstyp für ihre Bereitstellung auszuwählen. Das folgende Beispiel setzt den Aktivierungstyp für MoE-Schichten auf MXFP8:

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

Zukünftige Pläne

Breitere und tiefere Geräteunterstützung

Humming wird die Hardwareabdeckung weiter ausbauen und sich dabei an die Rechen- und Speicherfähigkeiten neuer Geräte anpassen, damit mehr Plattformen das Leistungspotenzial quantisierter Inferenz ausschöpfen können. Zu den Prioritäten gehören:

• Unterstützung weiterer Funktionen der Rubin-Architektur, einschließlich Locality Domains, um Task-Zuordnung und Datenzugriff weiter zu optimieren.
• Ausweitung der Unterstützung auf weitere T-Head-PPU-Geräte, einschließlich des Zhenwu M890.
• Unterstützung weiterer Geräte, die mit der CUDA-Programmierung kompatibel sind, um bestehende Quantisierungsformate, Kernel-Komponenten und Tuning-Fähigkeiten auf zusätzliche Hardwareplattformen zu bringen.

Hinzufügen von Autotune

Mit gezieltem Tuning kann Humming bereits in den meisten Szenarien State-of-the-Art-Leistung erreichen. Die große Vielfalt an Modellformen, Quantisierungskonfigurationen und Geräteressourcen macht es jedoch schwierig, dass Heuristiken in jedem Fall die optimale Konfiguration auswählen.

Humming plant daher die Einführung von Autotune, bei dem Kandidatenkonfigurationen auf dem Zielgerät benchmarket werden, um für jede Workload bessere Ausführungsparameter zu finden. Autotune wird bestehende Heuristiken ergänzen, die manuelle Konfigurationssuche reduzieren und weiteren Workloads helfen, sich der bestmöglichen Leistung der Kernel anzunähern.

Tiefere Integration mit Inferenz-Frameworks

Humming wird seine Integration mit Inferenz-Frameworks wie vLLM weiter vertiefen und die Unterstützung in den Bereichen Laden quantisierter Modelle, Backend-Auswahl, parallele Inferenz, Laufzeit-Scheduling und Geräteanpassung ausbauen. Wir wollen mehr Quantisierungsmethoden mit dem gesamten Umfang der Framework-Funktionen kompatibel machen, einschließlich Optimierungen wie quantisierter All-to-All-Kommunikation (A2A) für Expert Parallelism (EP). Diese Bemühungen werden leistungsstarke quantisierte Kernel in vielfältigen Bereitstellungsszenarien leichter nutzbar machen und dazu beitragen, Kernel-Leistungsgewinne in End-to-End-Inferenzvorteile umzusetzen.

Danksagungen

Wir danken Jinzhen Lin, Zekai Gu und Zhilin Huang von Ant Group sowie Michael Goin, Roberto L. Castro, Chauncey Jiang und weiteren Mitgliedern der vLLM-Community für ihre Beiträge und Unterstützung während der Entwicklung von Humming.

Wir danken auch der vLLM-Community für die Bereitstellung von Entwicklungs-, Test- und CI-Ressourcen, die es ermöglichten, Humming über ein breites Spektrum an Hardware und Workloads zu validieren und zu verbessern.

Hummings Design und Optimierung orientieren sich an Marlin, vLLM Marlin, LMDeploy, DeepGEMM, FlashInfer, CUTLASS und anderen Open-Source-Projekten. Wir sind ihren Autoren und Beitragenden dankbar, dass sie ihre Ideen und Implementierungen geteilt haben und so die leistungsstarke quantisierte Inferenz für die gesamte Community voranbringen.