DeepSeek dévoile DeepEP, une bibliothèque open source pour l’entraînement et l’inférence MoE

 (github.com/deepseek-ai)
1 points par GN⁺ 2025-02-26 | 1 commentaires | Partager sur WhatsApp
  • DeepEP est une bibliothèque de communication haute performance axée sur le parallélisme d’experts (EP) pour l’entraînement et l’inférence ML modernes, avec des kernels GPU all-to-all pour le dispatch/combine MoE et la prise en charge des basses précisions comme le FP8
  • La version V2 refactorise entièrement l’EP pour offrir des performances équivalentes ou supérieures à V1 avec beaucoup moins de ressources SM, et remplace le backend NVSHMEM par un NCCL Gin plus léger
  • Les tests, alignés sur une configuration V3, ont été réalisés avec 8K tokens par batch, hidden 7168, top 8 experts, dispatch FP8 et combine BF16 ; V2 atteint jusqu’à 1,3x de performance de pointe et jusqu’à 4x de réduction des SM par rapport à V1
  • Tous les kernels sont compilés à l’exécution sous forme de modules JIT légers, sans compilation CUDA à l’installation, et V2 unifie les API haut débit et faible latence via une interface unique ElasticBuffer
  • Un GPU Hopper SM90, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink et un réseau RDMA inter-nœuds sont requis ; Engram, PP et CP restent des fonctionnalités expérimentales

Ce que couvre DeepEP

  • DeepEP (DeepEveryParallel) est une bibliothèque de communication haute performance pour l’entraînement et l’inférence en machine learning modernes
  • Sa fonction centrale actuelle est le parallélisme d’experts (Expert Parallelism, EP), avec des kernels GPU all-to-all à haut débit et faible latence pour le dispatch et le combine des MoE
  • Elle prend en charge la communication en basse précision, y compris le FP8
  • Elle inclut aussi des primitives expérimentales pour le pipeline parallelism (PP), le context parallelism (CP) et l’accès mémoire distant (Engram)
  • Tous les kernels sont compilés à l’exécution sous forme de modules JIT (Just-In-Time) légers, sans nécessiter de compilation CUDA lors de l’installation
  • Malgré cette conception légère, l’objectif est d’atteindre, voire de dépasser, les limites de bande passante matérielle sur diverses configurations

Changements majeurs de la version V2

  • V2 est une version qui refactorise entièrement l’Expert Parallelism
    • Elle est conçue pour atteindre des performances extrêmes avec plusieurs fois moins de ressources SM que V1
    • Elle prend en charge des domaines de scale-up et scale-out plus vastes
    • Le backend passe de NVSHMEM à un backend NCCL Gin plus léger
  • Les nouvelles fonctionnalités incluent :
    • une compilation JIT complète
    • un backend NCCL Gin léger, uniquement en en-tête
    • la possibilité de réutiliser des communicateurs NCCL existants
    • dans EPv2, l’unification des API haut débit et faible latence dans une interface unique ElasticBuffer
    • un nouveau layout GEMM
    • la prise en charge de domaines de scale-up et scale-out plus grands, jusqu’à EP2048
    • le calcul analytique du nombre de SM et de QP, supprimant le besoin d’auto-tuning
    • le maintien de la prise en charge des modes hybride et direct
    • dans un entraînement legacy de type V3, une réduction de l’usage des SM de 24 à 4~6, avec des performances équivalentes ou meilleures
    • 0 SM Engram basé sur RDMA
    • 0 SM PP basé sur RDMA
    • 0 SM CP basé sur Copy Engine

Contraintes et fonctionnalités en cours

  • V2 consomme davantage de taille de buffer que V1
  • L’EP faible latence RDMA en 0 SM n’est plus pris en charge
  • Engram, PP et CP sont des fonctionnalités expérimentales
  • Les fonctionnalités en cours comprennent :
    • des Elastic GPU & CPU buffers, un espace d’adressage virtuel continu mappant de façon mixte la mémoire physique GPU et CPU
      • l’objectif est de permettre un Engram entièrement automatique et transparent, ou un EP déséquilibré
    • un travail sur la gestion du déséquilibre de charge via l’EP replay et la réduction de la taille des buffers intermédiaires
    • l’implémentation d’all-gather updates et de reduce-scatter pour DP et TP
  • La documentation de V1 basée sur NVSHMEM est disponible dans docs/legacy.md

Résultats de performance

  • Les tests ont été réalisés selon une configuration V3, avec les paramètres suivants :
    • 8K tokens par batch
    • dimension hidden 7168
    • top 8 experts
    • dispatch FP8
    • combine BF16
  • Principaux résultats :
    • 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
  • Les valeurs mesurées sont une bande passante logique ; par exemple, les 90 GB/s de EP 8 x 2 incluent le trafic de local rank
  • V2 atteint jusqu’à 1,3x de performance de pointe et réduit jusqu’à 4x le nombre de SM par rapport à V1
  • Les résultats pour des configurations EP plus grandes sont actuellement omis, et il est recommandé aux utilisateurs d’effectuer leurs propres benchmarks
  • D’après l’expérience interne, les kernels devraient continuer à saturer la bande passante matérielle même à plus grande échelle
  • Les données de performance de V1 se trouvent dans docs/legacy.md

Installation et prérequis

  • Les prérequis sont les suivants :
    • un GPU Hopper SM90 ou une architecture prenant en charge l’ISA PTX SM90
    • Python 3.8 ou plus
    • CUDA 12.3 ou plus pour GPU SM90
    • PyTorch 2.10 ou plus
    • NCCL 2.30.4 ou plus
    • NVLink pour la communication intra-nœud
    • un réseau RDMA pour la communication inter-nœuds
  • Il est recommandé d’installer NCCL via pip pour que DeepEP puisse le détecter automatiquement dans l’environnement Python
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • Pour la prise en charge des méthodes legacy, la bibliothèque dépend aussi de NVSHMEM ; voir les instructions dans le NVSHMEM Installation Guide
  • Exemples de build et d’exécution des tests en développement :
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
  • L’installation s’effectue avec la commande suivante :
python setup.py install
  • Après installation, deep_ep peut être importé et utilisé dans un projet Python

Interface centrée sur ElasticBuffer

  • Dans V2, toutes les opérations EP sont unifiées sous une interface unique ElasticBuffer
    • les API haut débit et faible latence y sont traitées de la même manière
    • le buffer peut être initialisé en spécifiant directement la configuration MoE
    • le nombre optimal de SM et de QP est calculé analytiquement
  • Un exemple d’initialisation de buffer utilise ElasticBuffer.get_buffer_size_hint() pour calculer la taille requise et vérifier si un buffer existant peut être réutilisé
  • Lors de la création d’un nouveau buffer, on spécifie notamment num_max_tokens_per_rank, hidden, num_topk, use_fp8_dispatch, etc.
  • _buffer.get_theoretical_num_sms(num_experts, num_topk) permet d’obtenir le nombre théorique de SM à utiliser pour le kernel de communication
  • Dans les appels dispatch et combine, il est possible de définir directement num_sms afin d’écraser la valeur calculée

Schémas d’usage pour l’entraînement, le prefill et le décodage

  • En entraînement ou en prefill d’inférence, le MoE dispatch route les tokens vers l’expert correspondant sur tous les ranks
    • il prend en charge les entrées BF16 et FP8
    • handle contient les métadonnées de routage nécessaires à l’appel combine ultérieur
    • handle.num_recv_tokens_per_expert_list fournit le nombre de tokens par expert nécessaire au GEMM
  • Le backward pass du MoE dispatch est en pratique géré par combine
  • Le MoE combine réduit les sorties des experts vers leur rank d’origine
  • Le backward pass du MoE combine est en pratique géré par dispatch
  • Le chevauchement entre communication et calcul est géré via l’interface EventOverlap
    • il permet d’exécuter des calculs indépendants pendant que la communication est en cours
    • avant d’utiliser le résultat, il faut synchroniser le compute stream avec event.current_stream_wait()
  • Le même ElasticBuffer est aussi utilisé pour le décodage en inférence
    • quand la décision de gating ne change pas, les métadonnées de routage peuvent être réutilisées via cached_handle
    • ce schéma évite le recalcul du layout et la synchronisation CPU

Variables d’environnement et constantes figées à la compilation

  • Paramètres généraux
    • EP_BUFFER_DEBUG : affiche les informations de débogage sur l’initialisation du buffer, l’approximation des SM et le backend
    • EP_SUPPRESS_NCCL_CHECK : supprime la vérification des incompatibilités de version NCCL
    • EP_AVOID_RECORD_STREAM : évite record_stream sur les tenseurs de sortie
    • EP_NUM_TOPK_IDX_BITS : redéfinit le nombre de bits pour l’encodage des index top-k
  • Paramètres réseau
    • EP_NIC_NAME : nom de NIC par défaut utilisé pour interroger les propriétés, valeur par défaut mlx5_0
    • EP_OVERRIDE_RDMA_SL : redéfinit l’index service level RDMA
    • EP_DISABLE_GIN : désactive le backend NCCL Gin
  • Paramètres JIT
    • EP_JIT_CACHE_DIR : répertoire de cache des kernels compilés, par défaut $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER : chemin du compilateur NVCC
    • EP_JIT_CPP_STANDARD : version du standard C++, par défaut 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM : paramètres liés au dump des sorties PTX/SASS
  • Certaines variables d’environnement ont un comportement persistent
    • elles sont capturées à la compilation et intégrées comme valeurs par défaut dans le package installé
    • ces valeurs par défaut s’appliquent automatiquement à l’import si elles ne sont pas redéfinies par les variables d’environnement courantes
    • les variables concernées sont EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR
  • Voir test_ep.py ou la documentation Python pour plus de détails

Recommandations de configuration réseau

  • DeepEP a été entièrement testé sur des réseaux InfiniBand
  • En théorie, il est aussi compatible avec RoCE (RDMA over Converged Ethernet)
  • Isolation du trafic
    • elle est prise en charge via les Virtual Lanes d’InfiniBand
    • il est recommandé de séparer les workloads expert-parallel des autres workloads sur des virtual lanes différentes
    • dans V2, l’affectation des virtual lanes peut être contrôlée via l’argument sl_idx ou la variable d’environnement EP_OVERRIDE_RDMA_SL
  • Adaptive routing
    • il s’agit d’une fonction de routage avancée des switches InfiniBand qui répartit uniformément le trafic sur plusieurs chemins
    • son activation est recommandée sous toutes les conditions de charge réseau, même si cela ajoute un peu de latence
  • Congestion control
    • il est recommandé de le désactiver, car il nuit à la bande passante maximale
    • si la congestion est inévitable, il est recommandé d’affecter le workload concerné à une virtual lane de plus faible priorité
  • PCI atomic mode
    • si le matériel le prend en charge, il est recommandé de configurer PCI_ATOMIC_MODE sur la NIC afin d’améliorer les performances des opérations atomiques RDMA
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Branches expérimentales et forks communautaires

  • Branches expérimentales
    • Zero-copy : supprime les copies entre les tenseurs PyTorch et les buffers de communication, réduisant fortement l’usage des SM pour les kernels généraux
    • Eager : utilise un protocole faible latence pour éliminer l’extra RTT latency ajoutée par les opérations atomiques RDMA
    • Hybrid-EP : nouvelle implémentation backend utilisant les instructions TMA, avec usage minimal des SM, prise en charge de domaines NVLink plus grands, chevauchement fin communication/calcul en single-batch, kernels PCIe et prise en charge de NVFP4
    • AntGroup-Opt : série d’optimisations réalisée par l’AntGroup Network Platform Department
    • Mori-EP : prise en charge du mode faible latence ROCm/AMD GPU basée sur le backend MORI
    • nvDev : branche basée sur V2 intégrant des fonctionnalités CUDA récentes comme Compute Fabric Transport
  • Forks communautaires
    • uccl/uccl-ep : prise en charge de l’exécution de DeepEP sur des GPU hétérogènes comme Nvidia et AMD, et des NIC comme EFA, Broadcom et CX7
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP : ajoute une solution multi-QP et la prise en charge des NIC double port pour le transport IBRC
    • antgroup/DeepXTrace : analyseur de diagnostic permettant de trouver efficacement et précisément les slow ranks
    • ROCm/mori : bibliothèque de communication AMD de nouvelle génération pour les workloads IA sensibles aux performances, incluant Wide EP, transfert de KVCache et collectives

Licence et citation

  • DeepEP V2 est construit sur le backend Gin de NCCL
  • Le code du dépôt est publié sous licence MIT
  • La référence de citation est DeepEP: an efficient expert-parallel communication library, avec l’année 2025

À lire aussi

1 commentaires

 
GN⁺ 2025-02-26
Avis sur Hacker News

  • Ils ont trouvé et utilisé une instruction PTX non documentée, ld.global.nc.L1::no_allocate.L2::256B, pour obtenir des performances extrêmes.
    Cette instruction accède à la mémoire GPU volatile avec le modificateur PTX en lecture seule non cohérente .nc, ce qui peut entraîner un comportement indéfini.
    Cela dit, sur l’architecture Hopper, il aurait été testé qu’elle garantit l’exactitude lorsqu’elle est utilisée avec .L1::no_allocate, tout en offrant de bien meilleures performances.

    • Concrètement, NVIDIA pourrait-elle plus tard, sur une nouvelle architecture, modifier subtilement le comportement de cette instruction hors documentation, intentionnellement ou non, et tout faire basculer ?

  • On a l’impression d’être un enfant dans une confiserie.
    Il y a beaucoup d’astuces qu’il faudrait trop longtemps à rétro-ingénierer correctement à partir des seuls articles, et j’espère que les publications de cette semaine ouvriront une renaissance où les MoE deviennent les modèles académiques par défaut.

    • Vu sous cet angle, je ne comprends pas ce qui se passe entre la pratique réelle des modèles de pointe et les modèles académiques.
      Depuis GPT-4, les premiers sont déjà tous des MoE, alors que les modèles publiés sont souvent des modèles denses, à l’exception de DeepSeek V3 et Mixtral.

  • On ne peut vraiment qu’aimer cette équipe.
    Elle repousse les limites de l’open source pour tout le monde.

    • Comme Open AI™, écrit avec une espace.
    • En réalité, ce n’est pas vraiment open source.
      Pour voir un vrai modèle open source, regardez OLMo 2 d’AI2 : https://allenai.org/blog/olmo2
      Ils partagent réellement tout ce qu’il faut pour reproduire le modèle, y compris les données elles-mêmes.
      Le lien ci-dessus dit aussi : « Comme la science ouverte complète exige plus que des poids publics, nous sommes heureux de partager avec la communauté plus large de la modélisation du langage une nouvelle mise à jour d’OLMo comprenant les poids, les données, le code, les recettes, les checkpoints intermédiaires et des modèles affinés par instructions. »

  • Zuckerberg devrait arrêter de prétendre que Meta publie son IA en open source.
    Ils diffusent même des publicités télévisées, mais en réalité ils ne publient que les poids, sans code.
    La seule vraie IA open source, c’est DeepSeek.

    • À strictement parler, DeepSeek n’est pas aussi open source qu’OLMo ou Open Euro.
      Parce qu’ils n’ont pas publié les données.
    • DeepSeek n’est clairement pas du vrai open source.
      Pour l’être, il faudrait utiliser une véritable licence open source, comme celles listées par l’OSI, et partager le code de pré-entraînement et de post-entraînement, le code lié au tuning, le code d’évaluation, tout ce qui concerne la sûreté et la censure, et probablement l’ensemble des données d’entraînement.
      Sinon, il est impossible de reproduire les poids, et partager des poids revient un peu à partager un programme compilé.
      À ma connaissance, le seul vrai modèle open source compétitif est OLMo 2 d’AI2 : https://allenai.org/blog/olmo2
      Ils ont aussi récemment publié une app d’inférence sur appareil, elle aussi open source : https://allenai.org/blog/olmoe-app
      Il y a aussi un autre modèle, Tülu 3, qui serait plus performant que DeepSeek V3 : https://allenai.org/blog/tulu-3-405B
    • Meta peaufine PyTorch depuis plus de dix ans.
      Presque tout ce qu’il faut pour entraîner des LLM, y compris les techniques les plus récentes, s’y trouve.
      Que faut-il de plus ? Des morceaux de code spécifiques à l’infrastructure de Meta ?
    • PyTorch est-il inclus aussi ?
    • Poids publics = gros tas de binaire.
      On est revenus au modèle FREEWARE / SHAREWARE.
      C’est ce genre de vocabulaire qu’il faudrait utiliser pour les poids « publics ».

  • Elle fournit une communication all-to-all efficace et optimisée, avec prise en charge intra-nœud et inter-nœuds via NVLink et RDMA, des kernels à haut débit pour l’entraînement et le prefill d’inférence, des kernels à faible latence pour le décodage d’inférence, la prise en charge native du dispatch FP8, et un contrôle flexible des ressources GPU pour le recouvrement calcul-communication.
    X : https://x.com/deepseek_ai/status/1894211757604049133

  • Les motivations derrière le travail de DeepSeek sont peut-être mauvaises.
    Par exemple, cela pourrait être une tentative soutenue par l’État de réduire à zéro l’avantage des États-Unis en matière d’IA, mais l’effet net pour le reste du monde est tout simplement excellent.
    Dans le pire des cas, même si c’est fait pour de mauvaises raisons, merci à DeepSeek : ils font réellement ce qu’OpenAI a menti pendant des années en prétendant faire pour le monde entier.

    • Dans le domaine des relations internationales, le bien et le mal ne s’appliquent pas vraiment de façon aussi tranchée.
      Publier cela en open source est-il plus « mal » qu’une interdiction d’exportation de GPU Nvidia haut de gamme ?
      La publication open source de DeepSeek est probablement simplement une chose bénéfique à la fois pour le Parti communiste chinois, avec son accord, et pour la communauté open source IA au sens large ; il ne faut pas y voir une position de principe.
      Trouver des moyens d’annuler l’avantage compétitif d’autres pays est une activité majeure de tous les gouvernements, grands comme petits.

  • C’est le deuxième round de publications open source sous licence MIT par la véritable entreprise Open AI™.
    Une fois de plus, DeepSeek est plus ouvert que l’entreprise à 157 milliards de dollars qui prétend être « Open ».
    Presque personne ne parle du Llama de Meta, mais tout le monde devrait s’attendre à ce que Llama 4 arrive avec des capacités de raisonnement.
    L’objectif est de ne pas se faire écraser au milieu d’une course vers zéro.

  • Pendant que les États-Unis fouillent les reçus de GPU à Singapour pour vérifier si DeepSeek n’a utilisé que des H800, le reste du monde peut donc faire tourner ces optimisations sur de vrais H100 ?
    Tout ça parce que les sanctions américaines rendaient les H100 difficiles à obtenir ou à utiliser, tout en continuant à faire comme si, par arrogance, les États-Unis croyaient que leurs ordres s’appliquaient au monde entier ?
    C’est bien comme ça qu’il faut le comprendre ?

  • Le PTX que tout le monde attendait est-il inclus cette fois ?

    • Oui, il y en a une partie dans le répertoire csrc/kernels.
      Cherchez asm pour trouver où il est utilisé.
    • Il faudrait expliquer au reste d’entre nous pourquoi le PTX que tout le monde attendait est si important.

  • Ça me rappelle les années 80-90, quand les gens bidouillaient l’assembleur ou cherchaient des instructions non documentées pour tirer le maximum des CPU.
    Un jour, les compilateurs seront probablement suffisamment optimisés, ou les GPU tellement puissants, pour que ce genre d’astuces ne fasse plus une grande différence, comme avec les CPU aujourd’hui.