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

 (github.com/deepseek-ai)
1 Punkte von GN⁺ 2025-02-26 | 1 Kommentare | Auf WhatsApp teilen
  • DeepEP ist eine High-Performance-Kommunikationsbibliothek mit Fokus auf Expert Parallelism (EP) in modernem ML-Training und -Inference. Sie bietet All-to-all-GPU-Kernel für MoE-Dispatch/Combine sowie Unterstützung für niedrige Präzision wie FP8
  • Das V2-Release refaktoriert EP grundlegend, erreicht mit deutlich weniger SM-Ressourcen als V1 gleiche oder bessere Leistung und wechselt das Backend von NVSHMEM zum leichteren NCCL Gin
  • In Tests mit V3-Konfiguration wurde mit 8K tokens per batch, hidden 7168, top 8 experts, FP8 dispatch und BF16 combine gemessen; V2 erzielte gegenüber V1 bis zu 1,3-fache Spitzenleistung und bis zu 4-fache SM-Einsparung
  • Alle Kernel werden als leichtgewichtiges JIT-Modul zur Laufzeit kompiliert, bei der Installation ist keine CUDA-Kompilierung nötig, und V2 vereinheitlicht High-Throughput- und Low-Latency-APIs in einer einzigen ElasticBuffer-Schnittstelle
  • Erforderlich sind Hopper-SM90-GPU, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink und ein RDMA-Netzwerk zwischen Knoten; Engram, PP und CP sind experimentelle Funktionen

Umfang von DeepEP

  • DeepEP (DeepEveryParallel) ist eine High-Performance-Kommunikationsbibliothek für modernes Machine-Learning-Training und -Inference
  • Die aktuelle Kernfunktion ist Expert Parallelism (EP); dafür stellt sie High-Throughput- und Low-Latency-All-to-all-GPU-Kernel für MoE-Dispatch und -Combine bereit
  • Unterstützt wird Low-Precision-Kommunikation, einschließlich FP8
  • Enthalten sind außerdem experimentelle Primitive für Pipeline Parallelism (PP), Context Parallelism (CP) und Remote Memory Access (Engram)
  • Alle Kernel werden als leichtgewichtige JIT(Just-In-Time)-Module zur Laufzeit kompiliert; im Installationsprozess ist keine CUDA-Kompilierung erforderlich
  • Trotz des leichtgewichtigen Designs zielt DeepEP in mehreren Konfigurationen auf eine Leistung nahe an oder über den Grenzen der Hardware-Bandbreite

Zentrale Änderungen im V2-Release

  • V2 ist ein Release mit vollständig refaktoriertem Expert Parallelism
    • Entwickelt, um extreme Leistung mit um ein Mehrfaches weniger SM-Ressourcen als V1 zu erreichen
    • Unterstützt größere Scale-up- und Scale-out-Domänen
    • Das Backend wurde von NVSHMEM auf das leichtere NCCL Gin backend umgestellt
  • Neue Funktionen sind:
    • Vollständige JIT-Kompilierung
    • Header-only, leichtgewichtiges NCCL Gin backend
    • Wiederverwendung bestehender NCCL communicator möglich
    • Vereinheitlichung der High-Throughput-API und Low-Latency-API in EPv2 unter einer einzigen ElasticBuffer-Schnittstelle
    • Neues GEMM-Layout
    • Unterstützung größerer Scale-up- und Scale-out-Domänen bis EP2048
    • Analytische Berechnung der Anzahl von SMs und QPs, sodass kein Auto-Tuning mehr nötig ist
    • Hybrid Mode und Direct Mode werden weiterhin unterstützt
    • In V3-ähnlichem Legacy-Training wird die SM-Nutzung von 24 auf 4–6 reduziert, bei gleicher oder besserer Leistung
    • RDMA-basiertes 0 SM Engram
    • RDMA-basiertes 0 SM PP
    • Copy-Engine-basiertes 0 SM CP

Einschränkungen und laufende Funktionen

  • V2 hat einen höheren Buffer-Größenverbrauch als V1
  • 0-SM-RDMA-Low-Latency-EP wird nicht mehr unterstützt
  • Engram, PP und CP sind experimentelle Funktionen
  • Laufende Arbeiten umfassen:
    • Elastic GPU & CPU buffers, ein zusammenhängender virtueller Adressraum, der physischen GPU- und CPU-Speicher gemischt abbildet
      • Dies zielt darauf ab, vollständig automatisches und transparentes Engram oder unausgewogenes EP zu ermöglichen
    • Arbeiten, um per EP replay Lastungleichgewichte zu behandeln und die Größe temporärer Buffer zu reduzieren
    • All-gather-Updates und Reduce-scatter-Implementierung für DP und TP
  • Die NVSHMEM-basierte V1-Dokumentation ist unter docs/legacy.md verfügbar

Ergebnisse der Leistungsmessungen

  • Getestet wurde passend zur V3-Konfiguration unter folgenden Bedingungen:
    • 8K tokens pro Batch
    • hidden dimension 7168
    • top 8 experts
    • FP8 dispatch
    • BF16 combine
  • Die wichtigsten Ergebnisse:
    • SM90, CX7, EP 8 x 2: dispatch 90 GB/s RDMA, combine 81 GB/s RDMA, 12 SM
    • SM90, CX7, EP 8 x 4: dispatch 61 GB/s RDMA, combine 61 GB/s RDMA, 6 SM
    • SM100, CX7, EP 8 x 2: dispatch 90 GB/s RDMA, combine 91 GB/s RDMA, 12 SM
    • SM100, EP 8: dispatch 726 GB/s NVLink, combine 740 GB/s NVLink, 64 SM
    • SM100, EP 8: dispatch 643 GB/s NVLink, combine 675 GB/s NVLink, 24 SM
  • Die Messwerte sind logische Bandbreite; zum Beispiel enthält die Angabe von 90 GB/s bei EP 8 x 2 auch local rank traffic
  • V2 erreicht gegenüber V1 bis zu 1,3-fache Spitzenleistung und spart bis zu 4-fache SM-Anzahl ein
  • Ergebnisse für größere EP-Konfigurationen wurden derzeit ausgelassen; Nutzern wird empfohlen, eigene Benchmarks durchzuführen
  • Nach interner Erfahrung wird erwartet, dass die Kernel auch bei größerer Skalierung die Hardware-Bandbreite weiterhin sättigen
  • V1-Leistungsdaten finden sich in docs/legacy.md

Installation und Anforderungen

  • Die Anforderungen sind:
    • Hopper-SM90-GPU oder Architektur mit Unterstützung für SM90 PTX ISA
    • Python 3.8 oder höher
    • CUDA 12.3 oder höher für SM90-GPUs
    • PyTorch 2.10 oder höher
    • NCCL 2.30.4 oder höher
    • NVLink für Kommunikation innerhalb eines Knotens
    • RDMA-Netzwerk für Kommunikation zwischen Knoten
  • Es wird empfohlen, NCCL per pip zu installieren, damit DeepEP es in der Python-Umgebung automatisch erkennen kann
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • Für die Unterstützung von Legacy-Methoden hängt DeepEP außerdem von NVSHMEM ab; Installationshinweise finden sich im NVSHMEM Installation Guide
  • Beispiele für Build und Testausführung während der Entwicklung:
python setup.py build
ln -s build/lib.linux-x86_64-cpython-38/deep_ep_cpp.cpython-38-x86_64-linux-gnu.so

python tests/elastic/test_ep.py
python tests/elastic/test_agrs.py
python tests/elastic/test_engram.py
python tests/elastic/test_pp.py
  • Die Installation erfolgt mit folgendem Befehl:
python setup.py install
  • Nach der Installation kann deep_ep in Python-Projekten importiert und genutzt werden

Schnittstelle rund um ElasticBuffer

  • In V2 sind alle EP-Operationen unter einer einzigen ElasticBuffer-Schnittstelle zusammengeführt
    • High-Throughput- und Low-Latency-APIs werden über dieselbe Schnittstelle behandelt
    • Buffer können durch direkte Angabe der MoE-Konfiguration initialisiert werden
    • Die optimale Anzahl von SMs und QPs wird analytisch berechnet
  • Ein Beispiel für die Buffer-Initialisierung berechnet die benötigte Größe mit ElasticBuffer.get_buffer_size_hint() und prüft, ob ein vorhandener Buffer wiederverwendet werden kann
  • Beim Erstellen eines neuen Buffers werden unter anderem num_max_tokens_per_rank, hidden, num_topk und use_fp8_dispatch angegeben
  • Mit _buffer.get_theoretical_num_sms(num_experts, num_topk) erhält man die theoretische Anzahl der für den Kommunikationskernel zu verwendenden SMs
  • Wenn bei dispatch- und combine-Aufrufen num_sms direkt angegeben wird, überschreibt dies den berechneten Wert

Nutzungsmuster für Training, Prefill und Decoding

  • Beim Training oder Inference-Prefill routet MoE dispatch Tokens an den jeweiligen Expert auf allen Ranks
    • Unterstützt werden BF16- und FP8-Eingaben
    • handle enthält die Routing-Metadaten, die für den späteren Combine-Aufruf benötigt werden
    • handle.num_recv_tokens_per_expert_list liefert die für GEMM benötigte Token-Anzahl pro Expert
  • Der Backward Pass von MoE dispatch wird tatsächlich mit combine verarbeitet
  • MoE combine reduziert die Expert-Ausgaben zurück auf den ursprünglichen Rank
  • Der Backward Pass von MoE combine wird tatsächlich mit dispatch verarbeitet
  • Die Überlappung von Kommunikation und Berechnung wird über die EventOverlap-Schnittstelle verwaltet
    • Während die Kommunikation läuft, kann unabhängige Berechnung ausgeführt werden
    • Vor der Nutzung des Ergebnisses wird der Compute Stream mit event.current_stream_wait() synchronisiert
  • Auch beim Inference-Decoding wird derselbe ElasticBuffer genutzt
    • Wenn sich die Gating-Entscheidung nicht ändert, werden Routing-Metadaten mit cached_handle wiederverwendet
    • Dieses Muster vermeidet Layout-Neuberechnung und CPU-Synchronisierung

Umgebungsvariablen und zur Build-Zeit festgelegte Werte

  • Allgemeine Einstellungen
    • EP_BUFFER_DEBUG: gibt Informationen zu Buffer-Initialisierung, SM-Approximation und Backend-Debugging aus
    • EP_SUPPRESS_NCCL_CHECK: unterdrückt die Prüfung auf NCCL-Versionsinkompatibilität
    • EP_AVOID_RECORD_STREAM: vermeidet record_stream für Ausgabe-Tensoren
    • EP_NUM_TOPK_IDX_BITS: überschreibt die Anzahl der Bits für die top-k-index-Codierung
  • Netzwerkeinstellungen
    • EP_NIC_NAME: Standard-NIC-Name für die Abfrage von NIC-Eigenschaften, Standardwert mlx5_0
    • EP_OVERRIDE_RDMA_SL: überschreibt den RDMA service level index
    • EP_DISABLE_GIN: deaktiviert das NCCL Gin backend
  • JIT-Einstellungen
    • EP_JIT_CACHE_DIR: Cache-Verzeichnis für kompilierte Kernel, Standardwert $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: Pfad zum NVCC-Compiler
    • EP_JIT_CPP_STANDARD: C++-Standardversion, Standardwert 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: Einstellungen zum Dump von PTX-/SASS-Ausgaben
  • Einige Umgebungsvariablen verhalten sich persistent
    • Sie werden zur Build-Zeit erfasst und als Standardwerte in das Installationspaket aufgenommen
    • Wenn sie beim Import nicht durch aktuelle Umgebungsvariablen überschrieben werden, werden diese Standardwerte automatisch angewendet
    • Die betroffenen Variablen sind EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS und EP_NCCL_ROOT_DIR
  • Weitere Details finden sich in test_ep.py oder in der Python-Dokumentation

Empfehlungen zur Netzwerkkonfiguration

  • DeepEP wurde vollständig in InfiniBand-Netzwerken getestet
  • Theoretisch ist es auch mit RDMA over Converged Ethernet, also RoCE, kompatibel
  • Traffic-Isolierung
    • Wird über Virtual Lanes von InfiniBand unterstützt
    • Es wird empfohlen, Expert-Parallel-Workloads und andere Workloads auf unterschiedliche Virtual Lanes zu trennen
    • In V2 kann die Zuweisung von Virtual Lanes über das Argument sl_idx oder die Umgebungsvariable EP_OVERRIDE_RDMA_SL gesteuert werden
  • Adaptive Routing
    • Eine erweiterte Routing-Funktion, bei der InfiniBand-Switches Traffic gleichmäßig über mehrere Pfade verteilen
    • Die Aktivierung wird für alle Netzwerklastbedingungen empfohlen, auch wenn zusätzliche Latenz entsteht
  • Congestion Control
    • Sollte deaktiviert werden, da sie der maximalen Bandbreite schadet
    • Wenn Überlastung unvermeidbar ist, wird empfohlen, den betreffenden Workload einer Virtual Lane mit niedrigerer Priorität zuzuweisen
  • PCI atomic mode
    • Wenn die Hardware dies unterstützt, wird empfohlen, PCI_ATOMIC_MODE der NIC zu setzen, um die Leistung von RDMA atomic operation zu verbessern
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Experimentelle Branches und Community-Forks

  • Experimentelle Branches
    • Zero-copy: Entfernt Kopien zwischen PyTorch-Tensoren und Kommunikations-Buffer und reduziert so die SM-Nutzung normaler Kernel deutlich
    • Eager: Nutzt ein Low-Latency-Protokoll, um die extra RTT latency durch RDMA atomic OPs zu entfernen
    • Hybrid-EP: Implementierung eines neuen Backends mit TMA instructions; umfasst minimale SM-Nutzung, Unterstützung größerer NVLink-Domänen, feingranulare Überlappung von Kommunikation und Berechnung für Single-Batch, PCIe-Kernel und NVFP4-Unterstützung
    • AntGroup-Opt: Eine Optimierungsserie des AntGroup Network Platform Department
    • Mori-EP: Unterstützung eines ROCm/AMD-GPU-Low-Latency-Modus auf Basis des MORI-Backends
    • nvDev: V2-basierter Branch mit aktuellen CUDA-Funktionen wie Compute Fabric Transport
  • Community-Forks
    • uccl/uccl-ep: Unterstützt die Ausführung von DeepEP auf heterogenen GPUs wie Nvidia und AMD sowie auf NICs wie EFA, Broadcom und CX7
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: Ergänzt IBRC transport um eine multi-QP-Lösung und Unterstützung für dual-port NICs
    • antgroup/DeepXTrace: Diagnostischer Analyzer, der slow ranks effizient und präzise findet
    • ROCm/mori: AMDs nächste Kommunikationsbibliothek für performancekritische AI-Workloads wie Wide EP, KVCache transfer und Collectives

Lizenz und Zitierung

  • DeepEP V2 baut auf dem NCCL Gin backend auf
  • Der Repository-Code ist unter der MIT License veröffentlicht
  • Der Zitier-Eintrag lautet DeepEP: an efficient expert-parallel communication library, das Jahr ist mit 2025 angegeben

Verwandte Beiträge

1 Kommentare

 
GN⁺ 2025-02-26
Hacker-News-Kommentare

  • Es wurde eine undokumentierte PTX-Instruktion ld.global.nc.L1::no_allocate.L2::256B gefunden und genutzt, um extreme Performance zu erzielen.
    Diese Instruktion greift mit dem nicht kohärenten, nur lesenden PTX-Modifikator .nc auf flüchtigen GPU-Speicher zu und kann daher undefiniertes Verhalten auslösen.
    Auf der Hopper-Architektur soll jedoch getestet worden sein, dass sie in Kombination mit .L1::no_allocate korrekt funktioniert und deutlich bessere Performance liefert.

    • Kann NVIDIA realistisch gesehen später auf einer neuen Architektur das Verhalten dieser nicht dokumentierten Instruktion subtil ändern und damit, ob beabsichtigt oder nicht, alles auf den Kopf stellen?

  • Es fühlt sich an wie ein Kind im Süßwarenladen.
    Es gibt viele Tricks, deren sauberes Reverse Engineering allein aus dem Paper viel zu lange dauern würde, und hoffentlich eröffnen die Veröffentlichungen dieser Woche eine Renaissance, in der MoE als akademisches Standardmodell genutzt wird.

    • Aus dieser Perspektive verstehe ich nicht, was zwischen tatsächlicher State-of-the-Art-Modellpraxis und akademischen Modellen passiert.
      Erstere sind seit GPT-4 praktisch alle MoE, während offene Modelle, abgesehen von DeepSeek V3 und Mixtral, oft weiterhin Dense Models sind.

  • Dieses Team muss man einfach mögen.
    Sie verschieben für alle die Grenzen von Open Source.

    • Im Sinne von Open AI™ mit Leerzeichen.
    • Es ist eigentlich nicht Open Source.
      Wer ein wirkliches Open-Source-Modell sehen will, sollte sich OLMo 2 von AI2 ansehen: https://allenai.org/blog/olmo2
      Sie teilen tatsächlich alles, was zur Reproduktion des Modells nötig ist, einschließlich der Daten selbst.
      Im obigen Link heißt es ebenfalls: „Da vollständig offene Wissenschaft mehr erfordert als öffentliche Gewichte, freuen wir uns, neue OLMo-Updates mit der breiteren Sprachmodellierungs-Community zu teilen, darunter Gewichte, Daten, Code, Rezepte, Zwischen-Checkpoints und Instruction-Tuning-Modelle.“

  • Zuckerberg sollte aufhören zu behaupten, Meta veröffentliche AI als Open Source.
    Sie schalten sogar TV-Werbung dafür, aber tatsächlich veröffentlichen sie nur Gewichte und keinen Code.
    Die einzige echte Open-Source-AI ist DeepSeek.

    • Streng genommen ist auch DeepSeek nicht so Open-Source-mäßig wie OLMo oder Open Euro.
      Denn sie haben die Daten nicht veröffentlicht.
    • DeepSeek ist eindeutig keine echte Open Source.
      Um Open Source zu sein, müsste es eine echte Open-Source-Lizenz verwenden, wie sie etwa die OSI aufführt, und Pre- und Post-Training-Code, Tuning-bezogenen Code, Evaluation-Code, alles zu Sicherheit und Zensur sowie wahrscheinlich den gesamten Trainingsdatensatz teilen.
      Andernfalls lassen sich die Gewichte nicht reproduzieren, und das Teilen von Gewichten ähnelt dem Teilen eines kompilierten Programms.
      Soweit ich weiß, ist das einzige wettbewerbsfähige echte Open-Source-Modell OLMo 2 von AI2: https://allenai.org/blog/olmo2
      Kürzlich haben sie auch eine App für On-Device-Inference veröffentlicht, die ebenfalls Open Source ist: https://allenai.org/blog/olmoe-app
      Außerdem gibt es ein weiteres Modell namens Tülu 3, das angeblich besser abschneidet als DeepSeek V3: https://allenai.org/blog/tulu-3-405B
    • Meta hat PyTorch über mehr als zehn Jahre hinweg verfeinert.
      Fast alles, was man zum Trainieren von LLMs einschließlich modernster Technik braucht, steckt darin.
      Was braucht man noch? Die Teile des Codes, die speziell auf Metas Infrastruktur zugeschnitten sind?
    • Zählt PyTorch auch?
    • Öffentliche Gewichte = Binär-Blob.
      Wir sind wieder beim FREEWARE-/SHAREWARE-Modell angekommen.
      Für „öffentliche“ Gewichte sollte man solche Begriffe verwenden.

  • Es bietet effiziente und optimierte All-to-all-Kommunikation, Unterstützung innerhalb eines Knotens und zwischen Knoten über NVLink und RDMA, High-Throughput-Kernels für Training und Inference-Prefilling, Low-Latency-Kernels für Inference-Decoding, native FP8-Dispatch-Unterstützung sowie flexible Steuerung von GPU-Ressourcen zur Überlappung von Rechen- und Kommunikationsarbeit.
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • Die Motivation hinter DeepSeeks Arbeit könnte falsch sein.
    Es könnte zum Beispiel ein staatlich unterstützter Versuch sein, den Vorsprung der USA bei AI auf null zu reduzieren, aber der Nettoeffekt für alle weltweit ist einfach großartig.
    Im schlimmsten Fall, also selbst wenn sie es aus den falschen Gründen tun, bin ich DeepSeek dankbar; sie tun tatsächlich das, was OpenAI der Welt über Jahre hinweg fälschlich versprochen hat.

    • Im Bereich internationaler Beziehungen gelten Richtig und Falsch nicht in diesem Maße.
      Ist es „falscher“, dies als Open Source zu veröffentlichen, als ein Exportverbot für hochwertige Nvidia-GPUs?
      DeepSeeks Open-Source-Veröffentlichung ist vermutlich einfach etwas, das mit Zustimmung der Kommunistischen Partei Chinas gleichzeitig der KPCh und der breiteren Open-Source-AI-Community nützt, und sollte nicht als prinzipielle Haltung verstanden werden.
      Wege zu finden, den Wettbewerbsvorteil anderer Länder zu beseitigen, gehört zu den Hauptaktivitäten aller großen und kleinen Regierungen.

  • Das ist Runde 2 der Open-Source-Veröffentlichungen des tatsächlichen Open AI™-Unternehmens unter MIT-Lizenz.
    Wieder einmal ist DeepSeek offener als ein 157-Milliarden-Dollar-Unternehmen, das von sich behauptet, „Open“ zu sein.
    Über Metas Llama spricht kaum jemand, aber alle sollten erwarten, dass Llama 4 mit Reasoning-Funktionen erscheint.
    Das Ziel ist, mitten im Wettlauf gegen null nicht zerquetscht zu werden.

  • Während die USA in Singapur GPU-Quittungen durchforsten, um zu prüfen, ob DeepSeek nur H800 verwendet hat, kann der Rest der Welt diese Optimierungen also auf vollwertigen H100 laufen lassen?
    Weil es wegen der US-Sanktionen schwierig war, H100 zu bekommen oder darauf zuzugreifen, und weil die USA in ihrer Arroganz weiterhin so tun, als würden ihre Anordnungen die ganze Welt erfassen?
    Verstehe ich das richtig?

  • Ist das PTX, auf das alle gewartet haben, diesmal enthalten?

    • Ja, einiges davon liegt im Verzeichnis csrc/kernels.
      Wenn man nach asm sucht, findet man die Stellen, an denen es verwendet wird.
    • Man müsste dem Rest von uns erklären, warum das PTX, auf das alle gewartet haben, so wichtig ist.

  • Das erinnert an die 80er und 90er, als Leute Assembly hackten oder undokumentierte Instruktionen suchten, um CPU-Performance herauszuholen.
    Irgendwann werden Compiler wohl ausreichend optimieren oder GPUs so leistungsfähig werden, dass solche Tricks, wie bei heutigen CPUs, keinen großen Unterschied mehr machen.