DeepSeek publie DeepEP, une bibliothèque open source pour l’entraînement et l’inférence MoE
(github.com/deepseek-ai)DeepEP
DeepEP est une bibliothèque de communication pour Mixture-of-Experts (MoE) et l’expert parallelism (EP). Elle fournit des kernels GPU all-to-all à haut débit et faible latence, connus sous le nom de dispatch et combine MoE. Elle prend également en charge les calculs en faible précision, y compris le FP8. Conformément à l’algorithme de group-limited gating proposé dans l’article DeepSeek-V3, elle fournit des kernels optimisés pour les transferts de bande passante inter-domaines asymétriques, en transférant les données du domaine NVLink vers le domaine RDMA. Ces kernels offrent un débit élevé, ce qui les rend adaptés à l’entraînement et aux tâches de préremplissage en inférence. Elle prend également en charge le contrôle du nombre de SM (Streaming Multiprocessors). Pour le décodage d’inférence sensible à la latence, DeepEP inclut des kernels à faible latence qui utilisent exclusivement RDMA afin de minimiser la latence. La bibliothèque introduit aussi une méthode de superposition communication-calcul basée sur des hooks, qui n’occupe pas les ressources SM.
Performances
Kernel général utilisant le transfert NVLink et RDMA
- Le kernel général a été testé sur H800 avec une bande passante NVLink maximale d’environ 160 GB/s, connecté à une carte réseau RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s de bande passante maximale).
- Suit la configuration de pré-entraînement DeepSeek-V3/R1 (4096 tokens par batch, hidden size 7168, top 4 groupes, top 8 experts, dispatch en FP8 et combine en BF16).
Kernel à faible latence utilisant uniquement RDMA
- Le kernel à faible latence a été testé sur H800 avec une carte réseau RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s de bande passante maximale).
- Suit une configuration de production typique de DeepSeek-V3/R1 (128 tokens par batch, hidden size 7168, top 8 experts, dispatch en FP8 et combine en BF16).
Démarrage rapide
Prérequis
- GPU Hopper (davantage d’architectures ou de dispositifs pourront être pris en charge ultérieurement)
- Python 3.8 ou version ultérieure
- CUDA 12.3 ou version ultérieure
- PyTorch 2.1 ou version ultérieure
- NVLink pour la communication intra-nœud
- Réseau RDMA pour la communication inter-nœuds
Télécharger et installer la dépendance NVSHMEM
DeepEP dépend d’une version modifiée de NVSHMEM. Il faut l’installer en suivant le guide d’installation.
Configuration réseau
DeepEP a été entièrement testé sur un réseau InfiniBand et est, en théorie, également compatible avec RDMA over Converged Ethernet (RoCE).
Isolation du trafic
InfiniBand prend en charge l’isolation du trafic via les Virtual Lanes (VL). Pour éviter les interférences entre différents types de trafic, il est recommandé de séparer les tâches sur des voies virtuelles comme suit :
- tâches utilisant le kernel général
- tâches utilisant le kernel à faible latence
- autres tâches
Dans DeepEP, l’affectation des voies virtuelles peut être contrôlée en définissant la variable d’environnement NVSHMEM_IB_SL.
Routage adaptatif
Le routage adaptatif est une fonctionnalité de routage avancée fournie par les switches InfiniBand, qui permet de répartir uniformément le trafic sur plusieurs chemins. À l’heure actuelle, le kernel à faible latence prend en charge le routage adaptatif, mais pas le kernel général (cela pourrait arriver bientôt). Activer le routage adaptatif pour le kernel général inter-nœuds peut provoquer des blocages ou des corruptions de données. Pour le kernel à faible latence, l’activation du routage adaptatif peut éliminer complètement la congestion réseau causée par les conflits de routage, mais ajoute une latence supplémentaire. Pour des performances optimales, la configuration suivante est recommandée :
- activer le routage adaptatif dans les environnements à forte charge réseau
- utiliser le routage statique dans les environnements à faible charge réseau
Contrôle de congestion
Le contrôle de congestion est désactivé, car aucune congestion significative n’a été observée en production.
Interface et exemples
Utilisation des exemples pour l’entraînement du modèle ou le préremplissage en inférence
Le kernel général peut être utilisé pour l’entraînement du modèle ou pendant l’étape de préremplissage en inférence (sans la partie rétropropagation).
Utilisation des exemples pour le décodage en inférence
Le kernel à faible latence peut être utilisé pendant l’étape de décodage en inférence.
Points d’attention
- Pour des performances extrêmes, ils ont découvert et utilisé l’instruction PTX non documentée
ld.global.nc.L1::no_allocate.L2::256B. Cette instruction provoque un comportement non défini lors de l’accès à la mémoire GPU volatile en utilisant un modificateur PTX non cohérent en lecture seule. Cependant, sur l’architecture Hopper, les tests avec.L1::no_allocateont montré des performances nettement supérieures. Si les kernels ne fonctionnent pas sur une autre plateforme, il est possible de désactiver cela en ajoutantDISABLE_AGGRESSIVE_PTX_INSTRS=1àsetup.py, ou d’ouvrir une issue. - Pour de meilleures performances sur un cluster, il est recommandé d’exécuter tous les tests et d’utiliser la configuration d’auto-tuning optimale. La configuration par défaut a été optimisée sur le cluster interne de DeepSeek.
Licence
Ce dépôt de code est publié sous licence MIT, et le code qui fait référence à NVSHMEM (y compris csrc/kernels/ibgda_device.cuh et third-party/nvshmem.patch) est soumis à la NVSHMEM SLA.
1 commentaires
Commentaires Hacker News
.L1::no_allocatesur l’architecture Hopper serait garantie, avec de bien meilleures performances