Accelerating GPU Inference of Large Language Models with Moderately Unstructured Sparse Weight Matrices

Technisches Resümee: Beschleunigung der GPU-Inferenz von Large Language Models mit moderat unstrukturierten spärlichen Gewichtmatrizen

Problemstellung
Die Bereitstellung von Large Language Models (LLMs) steht vor einem kritischen Engpass bei den Inferenzkosten. Während Techniken zur Modellpruning (z. B. Wanda, RIA) die Rechenlast durch die Einführung von Sparsity (Dünnbesetztheit) in Gewichtmatrizen reduzieren können, sind sie durch Anforderungen an die Modellqualität begrenzt. Im Gegensatz zu Modellen kleiner Skalierung reagieren LLMs hochsensibel auf hohe Sparsity-Grade; die Aufrechterhaltung der Performance beschränkt das Pruning typischerweise auf moderate unstrukturierte Sparsity (um 50 %).

Auf dieser spezifischen Sparsity-Stufe versagen bestehende GPU-Kernel für Sparse Matrix Multiplication (SpMM) darin, die Leistung von dichten Matrizenmultiplikations-Bibliotheken wie cuBLAS zu übertreffen. Diese Performance-Lücke resultiert aus drei grundlegenden Herausforderungen:

1. Inkompatibilität mit Tensor Cores: Moderne GPUs bieten hocheffiziente Sparse Tensor Cores, diese unterstützen jedoch strikt die strukturierte 2:4 Sparsity. Sie können irreguläre, unstrukturierte Pruning-Muster nicht direkt beschleunigen.
2. Metadaten-Overhead: Traditionelle Speicherformate (z. B. CSR) erfordern die Speicherung der vollständigen Positions-Metadaten für jedes Nicht-Null-Element. Bei einer 50 %igen Sparsity kann dieser Metadaten-Overhead so viel Platz beanspruchen wie die Werte selbst, was die Speicherbandbreiten-Vorteile zunichtemacht.
3. On-Chip-Decoding-Engpässe: Kompressionstechniken, die Metadaten reduzieren (z. B. Bitmap-Kodierung), verlagern die Arbeitslast auf das On-Chip-Decoding. Auf modernen GPUs mit High-Bandwidth Memory (HBM) kann der Zugriff auf den globalen Speicher nicht vollständig mit der Dekodierungsarbeit überlappen, die von langsameren CUDA-Cores ausgeführt wird, was zu einem Performance-Stall führt.

Methodik
Die Autoren schlagen eine ganzheitliche Lösung vor, die aus einem neuartigen dreischichtigen Matrizen-Speicherformat und einem kooptimierten SpMM-Kernel besteht, der für moderate unstrukturierte Sparsity entwickelt wurde.

1. Dreischichtiges Speicherformat
Das vorgeschlagene Format zerlegt die spärliche Matrix in drei distinkte Schichten, um ein Gleichgewicht zwischen Hardware-Kompatibilität und Speichereffizienz zu finden:

• Sparse-TC-Schicht: Diese Schicht extrahiert 2:4 strukturierte Sparsity-Muster aus der irregulären Verteilung. Sie behält zwei Nicht-Null-Werte von jeweils vier aufeinanderfolgenden Gewichten bei. Da das Muster fix ist, benötigt es nur 2 Bit pro Wert für die Metadaten, was eine direkte Kompatibilität mit Sparse Tensor Cores ohne Dekodierungs-Overhead gewährleistet.
• Slot-Filling-Schicht: Überschüssige Nicht-Null-Elemente, die nicht in das 2:4-Muster passen, werden in die durch die Sparse-TC-Schicht hinterlassenen freien Slots umverteilt. Um die Metadaten für diese ~17,7 % der Elemente zu minimieren, führen die Autoren eine Parallel Differential Distance (PDD) Kodierung ein. Anstatt absolute Indizes zu speichern, kodiert PDD die relativen Abstände zwischen den umverteilten Elementen. Die meisten Abstände sind klein genug, um in 4 Bit kodiert zu werden. Für größere Lücken werden Dummy-Punkte eingefügt, um die Spannweite zu brechen, sodass kein einzelner Abstand das Kodierungslimit überschreitet.
• Residual-Schicht: Die verbleibenden Nicht-Null-Elemente (typischerweise <1 % der Gesamtzahl), die nicht effizient durch die ersten beiden Schichten verarbeitet werden können, werden in einem traditionellen Compressed Sparse Row (CSR) Format gespeichert. Aufgrund des minimalen Datenvolumens ist der Metadaten-Overhead hier vernachlässigbar.

2. Kooptimierter SpMM-Kernel
Aufbauend auf diesem Format entwerfen die Autoren einen GPU-Kernel, der Sparse Tensor Cores und CUDA Cores gemeinsam in einem pipelinierten Ausführungsmodell nutzt:

• Paralleles Dekodierung: CUDA Cores sind mit der Dekodierung der PDD-Metadaten aus der Slot-Filling-Schicht beauftragt. Da diese Schicht nur einen kleinen Bruchteil der Gesamtdaten ausmacht, wird der Dekodierungs-Overhead effektiv kontrolliert.
• Strukturierte Beschleunigung: Sparse Tensor Cores verarbeiten die Sparse-TC-Schicht mittels wgmma-Instruktionen (Hopper-Architektur) und nutzen dabei die 2:4 Struktur für hohen Durchsatz.
• Residual-Verarbeitung: CUDA Cores verarbeiten auch die spärliche Residual-Schicht und schreiben die Ergebnisse in einen separaten On-Chip-Buffer, um sie am Ende der Pipeline mit den Tensor-Core-Outputs zu reduzieren.
• Pipeline-Überlappung: Der Kernel implementiert eine feingranulare Pipeline, welche drei kritische Operationen überlappt:
  • Laden dichter Eingangsmatrizen (im L2 gecached) und spärlicher Gewichtstiles (aus dem Global Memory).
  • Dekodieren der PDD-Metadaten auf den CUDA Cores.
  • Ausführen der Matrizenmultiplikation auf den Sparse Tensor Cores.
    Diese „Load-as-sparse, Compute-as-sparse“-Strategie stellt sicher, dass die On-Chip-Berechnung mit dem Zugriff auf den globalen Speicher überlappt, wodurch die HBM-Ausnutzung maximiert wird.

Wesentliche Beiträge

• Neuartiges Speicherformat: Ein dreischichtiges Format, das speziell auf moderate unstrukturierte Sparsity zugeschnitten ist und die Kompatibilität mit Sparse Tensor Cores ermöglicht, den Metadaten-Overhead durch PDD-Kodierung drastisch reduziert und die On-Chip-Dekodierungskosten minimiert.
• Kooptimierter Kernel: Ein SpMM-Kernel, der Sparse Tensor Cores und CUDA Cores effizient koordiniert und durch Pipelining die Überlappung von Speicherzugriff und Berechnung realisiert.
• Erster Sieg über Dense: Die Implementierung ist das erste Beispiel dafür, dass moderat unstrukturierte Sparse SpMM auf modernen GPUs mit HBM die dichte Matrizenmultiplikation übertreffen kann.

Ergebnisse
Die Evaluierungen wurden auf NVIDIA H100 SXM5 GPUs mit den Modellen OPT-30B und OPT-66B bei Sparsity-Leveln von 50 %, 60 % und 70 % durchgeführt.

• Kernel-Level Performance: Bei 50 % Sparsity erreicht die vorgeschlagene Methode eine bis zu 1,64× Beschleunigung gegenüber SpInfer (EuroSys'25 Best Paper) und übertrifft signifikant andere Baselines wie SparTA und FlashLLM. Sie übertrifft zudem die dichte cuBLAS-Performance.
• End-to-End Inferenz: Die Methode liefert eine bis zu 1,41× End-to-End Beschleunigung gegenüber FlashLLM (VLDB'24) und einen bis zu 41,4 % höheren Durchsatz im Vergleich zu FlashLLM auf OPT-30B bei einer Batch-Größe von 16.
• Speichereffizienz: Der Ansatz reduziert den Peak-Global-Memory-Verbrauch im Vergleich zu dichten Baselines bei 50 % Sparsity um etwa 21,4 %, weist jedoch aufgrund des Fehlens einer Bitmap-Kodierung für Metadaten einen etwas höheren Speicherverbrauch als SpInfer auf.
• Skalierbarkeit: Die Methode behält ihre Vorteile über die Sparsity-Level von 50 % bis 70 % hinweg bei.

Bedeutung und Ansprüche
Das Paper behauptet, eine kritische Lücke in der LLM-Inferenzoptimierung zu schließen: die Unfähigkeit bestehender Sparse-Kernel, moderate unstrukturierte Sparsity effektiv zu nutzen. Durch die Einführung eines Speicherformats, das die Lücke zwischen irregulären Pruning-Mustern und strukturierter Hardware-Beschleunigung schließt, zeigen die Autoren, dass Pruning tatsächliche Inferenzbeschleunigungen liefern kann, ohne die Modellqualität zu opfern.

Die Arbeit hebt hervor, dass eine hohe Sparsity (>70 %) die Modellqualität verschlechtert, während moderate Sparsity (~50 %) den „Sweet Spot“ für LLMs darstellt. Die vorgeschlagene Lösung erschließt das Performance-Potenzial dieses Sweet Spots, der zuvor aufgrund der Ineffizienzen bestehender Sparse-Kernel unzugänglich war. Die Autoren merken an, dass ihre Optimierung spezifisch auf die Decode-Phase der LLM-Inferenz abzielt (in der die Batch-Größen klein und speichergebunden sind), was sie hochgradig geeignet für den skalierbaren Einsatz in Produktionsumgebungen macht, in denen Prefill- und Decode-Phasen oft entkoppelt sind. Der Quellcode wird zur Verfügung gestellt, um die weitere Forschung und Adaption zu erleichtern.