DeepSeek veröffentlicht die Open-Source-Bibliothek DeepEP für MoE-Training und Inferenz

DeepEP

DeepEP ist eine Kommunikationsbibliothek für Mixture-of-Experts (MoE) und Expert Parallelism (EP). Sie bietet schnelle All-to-All-GPU-Kernel mit niedriger Latenz, bekannt als MoE-Dispatch und -Combine. Außerdem unterstützt sie Berechnungen mit niedriger Präzision einschließlich FP8. Passend zum im DeepSeek-V3-Paper vorgeschlagenen gruppenbegrenzten Gating-Algorithmus stellt sie Kernel bereit, die die asynchrone Domänenbandbreitenübertragung optimieren und Daten von der NVLink-Domäne zur RDMA-Domäne übertragen. Diese Kernel bieten hohen Durchsatz und eignen sich für Training sowie Prefill-Workloads bei der Inferenz. Außerdem wird die Steuerung der Anzahl der SMs (Streaming Multiprocessors) unterstützt. Für latenzsensitive Inferenz-Decodierung enthält DeepEP Low-Latency-Kernel, die reines RDMA nutzen, um die Latenz zu minimieren. Die Bibliothek führt zudem eine hook-basierte Methode zur Überlappung von Kommunikation und Berechnung ein, die keine SM-Ressourcen belegt.

Leistung

Allgemeine Kernel mit NVLink- und RDMA-Übertragung

• Auf H800 wurde der allgemeine Kernel mit maximal etwa 160 GB/s NVLink-Bandbreite sowie einer CX7-InfiniBand-400-Gb/s-RDMA-Netzwerkkarte (~50 GB/s maximale Bandbreite) getestet.
• Folgt dem DeepSeek-V3/R1-Pretraining-Setup (4096 Tokens pro Batch, 7168 Hidden, Top-4-Gruppen, Top-8-Experten, FP8-Dispatch und BF16-Combine).

Low-Latency-Kernel mit reinem RDMA

• Auf H800 wurde der Low-Latency-Kernel mit einer CX7-InfiniBand-400-Gb/s-RDMA-Netzwerkkarte (~50 GB/s maximale Bandbreite) getestet.
• Folgt dem typischen produktiven DeepSeek-V3/R1-Setup (128 Tokens pro Batch, 7168 Hidden, Top-8-Experten, FP8-Dispatch und BF16-Combine).

Schnellstart

Anforderungen

• Hopper-GPU (in Zukunft könnten weitere Architekturen oder Geräte unterstützt werden)
• Python 3.8 oder höher
• CUDA 12.3 oder höher
• PyTorch 2.1 oder höher
• NVLink für die Kommunikation innerhalb eines Knotens
• RDMA-Netzwerk für die Kommunikation zwischen Knoten

NVSHMEM-Abhängigkeit herunterladen und installieren

DeepEP hängt von einer modifizierten Version von NVSHMEM ab. Für die Installation sollte die Installationsanleitung konsultiert werden.

Netzwerkkonfiguration

DeepEP wurde vollständig in InfiniBand-Netzwerken getestet und ist theoretisch auch mit RDMA over Converged Ethernet (RoCE) kompatibel.

Traffic-Isolation

InfiniBand unterstützt Traffic-Isolation über Virtual Lanes (VL). Um Interferenzen zwischen unterschiedlichen Traffic-Typen zu vermeiden, wird empfohlen, Workloads wie folgt auf virtuelle Lanes aufzuteilen:

• Workloads mit allgemeinen Kerneln
• Workloads mit Low-Latency-Kerneln
• Sonstige Workloads

In DeepEP kann die Zuweisung virtueller Lanes durch Setzen der Umgebungsvariable NVSHMEM_IB_SL gesteuert werden.

Adaptives Routing

Adaptives Routing ist eine fortgeschrittene Routing-Funktion von InfiniBand-Switches, die Traffic gleichmäßig auf mehrere Pfade verteilen kann. Derzeit unterstützen die Low-Latency-Kernel adaptives Routing, die allgemeinen Kernel jedoch nicht (eine Unterstützung könnte bald folgen). Wird adaptives Routing für normale Inter-Node-Kernel aktiviert, kann es zu Deadlocks oder Problemen mit Datenkorruption kommen. Bei den Low-Latency-Kerneln kann die Aktivierung von adaptivem Routing Netzwerkstaus durch Routing-Konflikte vollständig beseitigen, verursacht jedoch zusätzliche Latenz. Für optimale Leistung wird folgende Konfiguration empfohlen:

• Adaptives Routing in Umgebungen mit hoher Netzwerkauslastung aktivieren
• Statisches Routing in Umgebungen mit geringer Netzwerkauslastung verwenden

Congestion Control

Da in Produktionsumgebungen keine signifikante Congestion beobachtet wurde, ist Congestion Control deaktiviert.

Schnittstellen und Beispiele

Beispielnutzung für Modelltraining oder Inferenz-Prefill

Die allgemeinen Kernel können für Modelltraining oder die Inferenz-Prefill-Phase (ohne Backward-Teil) verwendet werden.

Beispielnutzung für Inferenz-Decodierung

Die Low-Latency-Kernel können in der Inferenz-Decodierungsphase verwendet werden.

Hinweise

• Für extreme Performance wurde die undokumentierte PTX-Instruktion ld.global.nc.L1::no_allocate.L2::256B entdeckt und eingesetzt. Diese Instruktion führt undefiniertes Verhalten aus, da sie mit einem inkonsistenten schreibgeschützten PTX-Modifikator auf volatile GPU-Speicherbereiche zugreift. Tests mit .L1::no_allocate auf der Hopper-Architektur zeigten jedoch deutlich bessere Leistung. Falls die Kernel auf anderen Plattformen nicht funktionieren, kann dies deaktiviert werden, indem DISABLE_AGGRESSIVE_PTX_INSTRS=1 zu setup.py hinzugefügt wird, oder es kann ein Issue gemeldet werden.
• Für bessere Leistung im Cluster wird empfohlen, alle Tests auszuführen und die optimale Auto-Tuning-Konfiguration zu verwenden. Die Standardkonfiguration wurde für DeepSeeks internen Cluster optimiert.

Lizenz

Dieses Code-Repository wird unter der MIT-Lizenz veröffentlicht. Code, der auf NVSHMEM verweist (einschließlich csrc/kernels/ibgda_device.cuh und third-party/nvshmem.patch), unterliegt der NVSHMEM SLA.

[GN⁺]

Hacker-News-Kommentare

• Für extreme Performance wurden undokumentierte PTX-Instruktionen entdeckt und eingesetzt. Diese Instruktionen verwenden inkonsistente schreibgeschützte PTX-Modifikatoren für den Zugriff auf flüchtigen GPU-Speicher, was zu undefiniertem Verhalten führen könnte. Auf der Hopper-Architektur ist die Korrektheit mit .L1::no_allocate jedoch getestet und gewährleistet, und die Performance dürfte deutlich besser sein
• Zuckerberg sollte aufhören zu behaupten, Meta würde AI als Open Source veröffentlichen. Sie veröffentlichen nur Gewichte, nicht den Code. Echte Open-Source-AI gibt es nur von DeepSeek
• Ich fühle mich wie ein Kind im Süßwarenladen. Einige dieser Tricks würden viel zu lange brauchen, um sie auf Basis der Paper korrekt rückzuentwickeln. Hoffentlich leiten die Ankündigungen dieser Woche eine Renaissance von MoE als akademischem Standardmodell ein
• Man kann diese Leute einfach nur lieben. Sie verschieben die Grenzen von Open Source wirklich für uns alle. Danke fürs Teilen
  • effiziente und optimierte All-to-All-Kommunikation
  • Unterstützung innerhalb eines Knotens und zwischen Knoten über NVLink und RDMA
  • Kernel mit hohem Durchsatz für Training und Inference-Prefill
  • Kernel mit niedriger Latenz für Inference-Decoding
  • native Unterstützung für FP8-Dispatch
  • flexible Steuerung von GPU-Ressourcen zur Überlappung von Berechnung und Kommunikation
• Die Motivation hinter der Arbeit von DeepSeek mag fragwürdig sein (z. B. ein staatlich unterstützter Versuch, den technologischen Vorsprung der USA bei AI zu beseitigen). Aber weltweit ist das Ergebnis schlicht fantastisch
  • Selbst im schlimmsten Fall (wenn sie diese Arbeit aus den falschen Gründen machen) danke ich DeepSeek. Sie tun tatsächlich das, worüber OpenAI die ganze Welt seit Jahren belügt
  • Wirklich großartig
• Ich frage mich, ob das PTX, auf das alle gewartet haben, diesmal enthalten ist
• Die im technischen Bericht erwähnte PTX-Instruktion sollte hier mit dem Code verknüpft werden
• Während die USA in Singapur GPU-Belege nachverfolgen, um zu prüfen, ob DeepSeek nur H800 verwendet hat, kann der Rest der Welt diese Optimierungen auf vollständigen H100 ausführen
  • Ich frage mich, ob sie nur so tun, als wäre es wegen des Hochmuts der US-Sanktionen und ihres Glaubens, ihre Anordnungen würden die ganze Welt abdecken, schwierig gewesen, H100 zu bekommen oder darauf zuzugreifen
• Dies ist die zweite Open-Source-Veröffentlichung des eigentlichen Unternehmens „Open AI™“ und sie steht unter der MIT-Lizenz
  • DeepSeek ist offener als das Unternehmen, das behauptet, $157B+ wert zu sein
  • Fast niemand spricht über Metas Llama, und alle sollten erwarten, dass Llama 4 aus gutem Grund veröffentlicht wird
  • Das Ziel ist, im Rennen nach unten nicht in der Mitte stecken zu bleiben