- Bibliothèque de communication haute performance pour Mixture-of-Experts (MoE) et Expert Parallelism (EP)
- Fournit un kernel All-to-All basé sur GPU pour accélérer les opérations de dispatch et de combinaison des MoE
- Prend en charge les calculs en basse précision comme le FP8
- Applique l’algorithme de group-limited gating proposé dans l’article DeepSeek-V3 afin d’optimiser le forwarding de bande passante inter-domaines asymétriques
- Exemple : optimisation des transferts de données NVLink → RDMA
- Offre un débit élevé adapté à l’entraînement et aux tâches de prefilling d’inférence
- Inclut un kernel faible latence dédié à RDMA pour le décodage d’inférence sensible à la latence
- Fournit une technique de recouvrement communication-calcul (sans occuper les ressources SM)
Performances
Kernel général (transferts NVLink et RDMA)
- DeepEP a testé ses performances dans un environnement avec GPU H800 et réseau RDMA InfiniBand CX7 400 Gb/s
- Sur la base de la configuration DeepSeek-V3/R1, avec 4096 tokens par batch, 7168 nœuds cachés, une structure top-4 groupes et top-8 experts, en utilisant un dispatch FP8 et une combinaison BF16
- Les résultats montrent une bande passante de plus de 150 GB/s pour les communications intra-nœud (basées sur NVLink), et 40 à 47 GB/s pour les communications inter-nœuds (basées sur RDMA) selon le nombre d’experts
- Quand le nombre d’experts augmente, la bande passante RDMA tend à progresser légèrement (par exemple : 43 GB/s avec 16 experts, 46 GB/s avec 64 experts)
Kernel faible latence (RDMA pur)
- Les mesures du kernel faible latence montrent une latence nettement réduite par rapport au kernel général
- Dans un environnement traitant 128 tokens par batch, la latence augmente avec le nombre d’experts, tandis que la bande passante RDMA reste relativement stable
- Par exemple, elle passe de 163 microsecondes (us) avec 8 experts à 194 microsecondes (us) avec 256 experts
- Dans l’opération de combinaison (
combine), la latence est plus élevée que pour le dispatch, et la bande passante RDMA diminue progressivement sous les 40 GB/s à mesure que le nombre d’experts augmente
- En clair, le kernel faible latence est très rapide avec de petits groupes d’experts, mais quand leur nombre augmente, la latence s’allonge, ce qui impose de trouver un bon équilibre
Configuration réseau
Isolation du trafic (Traffic Isolation)
- Il est possible d’isoler le trafic en exploitant les Virtual Lanes (VL) d’InfiniBand
- Méthode de séparation recommandée :
- tâches utilisant le kernel général
- tâches utilisant le kernel faible latence
- autres tâches
- La configuration des VL peut se faire via la variable d’environnement
NVSHMEM_IB_SL
Routage adaptatif (Adaptive Routing)
- Prend en charge le routage adaptatif des switches InfiniBand
- Peut être activé avec le kernel faible latence, mais doit être désactivé avec le kernel général (risque de corruption des données s’il est activé)
- Recommandations de configuration :
- en cas de forte charge réseau : activer le routage adaptatif
- en cas de faible charge réseau : conserver un routage statique
Contrôle de congestion (Congestion Control)
- DeepEP fonctionne avec le contrôle de congestion désactivé
- Il a été confirmé que, dans un environnement réel, la congestion réseau n’est pas sévère
Principaux points techniques à considérer
- Utilisation d’instructions PTX non officielles :
ld.global.nc.L1::no_allocate.L2::256B est utilisé pour optimiser les performances
- Cela fonctionne correctement sur l’architecture Hopper, mais sur d’autres plateformes, il est possible de le désactiver avec
DISABLE_AGGRESSIVE_PTX_INSTRS=1
- Auto-tuning recommandé : pour obtenir les meilleures performances, il faut appliquer les réglages après des tests de performance spécifiques à chaque cluster
Aucun commentaire pour le moment.