Aller au contenu

7 · Multi-GPU et collectives

Quand un GPU ne suffit plus. Les quatre parallélismes, leurs coûts de communication, et la technique qui décide de tout : le recouvrement.

7.1 Les quatre parallélismes

Parallélisme Ce qu'on découpe Communication Usage
Données (DP) le lot all-reduce des gradients entraînement
Tenseurs (TP) les matrices de poids all-reduce par couche inférence, entraînement
Pipeline (PP) les couches send/recv entre étages très grands modèles
Experts (EP) les experts d'un MoE all-to-all par couche MoE
Séquence / contexte (SP/CP) la longueur de séquence all-gather / anneau contexte long

En pratique, on les combine : « 3D parallelism » (DP × TP × PP) est le standard pour l'entraînement de grands modèles, souvent étendu à 4D ou 5D avec EP et SP.

7.2 Le parallélisme de tenseurs, en détail

C'est le plus important pour l'inférence, et celui dont le coût est le plus directement lié aux noyaux.

Une couche linéaire \(\mathbf{y} = \mathbf{x}\mathbf{W}\) se découpe de deux façons.

Découpage en colonnes (\(\mathbf{W} = [\mathbf{W}_1 | \mathbf{W}_2]\)) :

\[\mathbf{y} = [\mathbf{x}\mathbf{W}_1 \,|\, \mathbf{x}\mathbf{W}_2]\]

Chaque GPU calcule une partie de la sortie. Pas de communication, mais la sortie est partitionnée.

Découpage en lignes (\(\mathbf{W} = [\mathbf{W}_1 ; \mathbf{W}_2]\), avec \(\mathbf{x} = [\mathbf{x}_1 | \mathbf{x}_2]\)) :

\[\mathbf{y} = \mathbf{x}_1\mathbf{W}_1 + \mathbf{x}_2\mathbf{W}_2\]

Chaque GPU calcule une somme partielle : il faut un all-reduce.

L'astuce standard, dans un bloc de transformeur : enchaîner colonnes puis lignes. QKV et gate/up sont découpés en colonnes, la projection de sortie et down en lignes. Résultat : deux all-reduce par couche au lieu de quatre.

Pour un modèle de 60 couches en TP8, cela fait 120 all-reduce par passe avant. Chacun échange \(b \times S \times d\) éléments.

Pourquoi le parallélisme de tenseurs ne passe pas l'échelle au-delà d'un nœud

Un all-reduce sur NVLink (900 Go/s bidirectionnels par GPU sur H100) est rapide. Sur InfiniBand entre nœuds (400 Gb/s = 50 Go/s), il est 18 fois plus lent.

Comme il y a deux all-reduce par couche et qu'ils sont sur le chemin critique, le TP au-delà d'un nœud est généralement rédhibitoire. C'est la raison pour laquelle les nœuds à 8 GPU avec NVSwitch sont l'unité de déploiement standard, et pourquoi on passe au pipeline ou aux experts au-delà.

7.3 NCCL

La bibliothèque de collectives de NVIDIA.

ncclAllReduce(sendbuff, recvbuff, count, ncclFloat, ncclSum, comm, stream);

Ce qu'elle fait sous le capot :

  • détecte la topologie : NVLink, NVSwitch, PCIe, InfiniBand, et leurs hiérarchies ;
  • choisit l'algorithme : anneau (bande passante optimale, latence en \(O(n)\)), arbre (latence en \(O(\log n)\)), ou hybrides ;
  • règle la taille des morceaux pour recouvrir communication et calcul de réduction ;
  • s'exécute sur un flux CUDA, donc concurremment avec vos noyaux.

Les collectives principales :

Opération Ce qu'elle fait
AllReduce somme (ou autre) sur tous, résultat partout
ReduceScatter somme, résultat partitionné
AllGather concatène les partitions, résultat partout
Broadcast un vers tous
AllToAll chacun envoie une part différente à chacun

Identité utile : \(\text{AllReduce} = \text{ReduceScatter} + \text{AllGather}\). C'est ce qu'exploite ZeRO/FSDP pour partitionner les états d'optimiseur.

7.4 NVSHMEM et la mémoire symétrique

NCCL a une limite : ses collectives sont des opérations à granularité de noyau. On ne peut pas commencer à communiquer les premiers résultats pendant que le noyau calcule les suivants.

NVSHMEM expose un espace d'adressage global partitionné (PGAS) directement au code du noyau :

__global__ void mon_noyau(float* local, float* symetrique, int pe_cible) {
    // calculer
    float resultat = calculer();

    // envoyer DEPUIS le noyau, sans revenir à l'hôte
    nvshmem_float_p(&symetrique[idx], resultat, pe_cible);

    // ou une écriture par blocs
    nvshmem_float_put(dest, source, nelems, pe_cible);

    nvshmem_quiet();               // s'assurer de la complétion
}

Chaque GPU alloue de la mémoire symétrique — une région à la même adresse virtuelle sur tous les processus — accessible par des opérations unilatérales (put, get, atomiques).

Ce que cela débloque :

  • envoyer les résultats au fur et à mesure de leur production ;
  • recouvrir communication et calcul à une granularité fine ;
  • implémenter des motifs de communication qui ne sont pas des collectives standard.

C'est la brique sur laquelle reposent DeepEP et les megakernels multi-GPU.

Les modes d'initiation

NVSHMEM permet les opérations initiées par l'hôte et initiées par le périphérique. Ce sont les secondes qui comptent ici : elles permettent la communication asynchrone à l'intérieur d'un noyau, et donc le recouvrement fin.

L'analyse système publiée en 2026 sur NVSHMEM (arXiv:2606.05951) détaille les compromis de performance entre les deux modes.

7.5 Le recouvrement calcul/communication

La technique qui décide de tout.

Sans recouvrement :

[calcul couche n][all-reduce][calcul couche n+1][all-reduce]…

Avec recouvrement :

[calcul couche n     ][calcul couche n+1     ]
        [all-reduce n         ][all-reduce n+1]

Les mécanismes :

Mécanisme Niveau
Flux CUDA séparés grossier, simple
Découpage en morceaux + collectives par morceau moyen
NVSHMEM depuis le noyau fin
Threads dédiés à la communication dans le noyau très fin

Le dernier est celui du megakernel de débit de Hazy Research : des threads storer dédiés gèrent la communication inter-GPU de façon asynchrone, via une abstraction de « Parallel Global Layout », ce qui libère les autres threads pour continuer à travailler au lieu de bloquer sur le réseau.

Un exemple : la transposition distribuée

Le même travail illustre une optimisation qui n'est pas exprimable avec des collectives standard.

Le problème : après l'attention en tensor-parallèle, il faut normalement un reduce-scatter après la projection de sortie.

Leur solution : répliquer la matrice de projection de sortie sur tous les GPU et effectuer cette projection en parallèle de données, ce qui supprime le reduce-scatter. Il est remplacé par une « transposition distribuée » qui repartitionne les données du format tensor-parallèle vers le format data-parallèle.

Résultat annoncé : le trafic réseau est divisé par huit. Et les auteurs notent que cette opération est « facilement exprimable dans le cadre du megakernel » mais « pas efficacement exprimable avec les motifs de communication standard ».

L'enseignement

Quand on contrôle la communication au niveau du noyau, on peut inventer des motifs qui n'existent pas dans NCCL. C'est un argument de fond en faveur des megakernels multi-GPU, indépendamment du gain sur le surcoût de lancement.

7.6 Les chiffres du multi-GPU

Bandes passantes, pour situer :

Lien Débit (par GPU)
HBM3 (H100) 3 350 Go/s
HBM3e (B200) 8 000 Go/s
NVLink 4 (H100) 900 Go/s bidirectionnels
NVLink 5 (B200) ~1 800 Go/s
PCIe 5.0 ×16 64 Go/s
InfiniBand NDR (400 Gb/s) 50 Go/s

La hiérarchie à retenir :

\[ \text{HBM} \gg \text{NVLink} \gg \text{PCIe} \approx \text{InfiniBand} \]

Un facteur ~4 entre HBM et NVLink, un facteur ~14 entre NVLink et le réseau. D'où la règle de conception : garder le trafic intensif à l'intérieur du nœud.

7.7 Le tableau de décision

Situation Parallélisme
Le modèle tient sur un GPU, entraînement DP (DDP ou FSDP)
Le modèle ne tient pas, ≤ 8 GPU TP dans le nœud
Le modèle ne tient pas, > 8 GPU TP dans le nœud + PP entre nœuds
MoE + EP
Contexte très long + SP/CP (Ring Attention, Ulysses)
Inférence latence-critique TP dans le nœud, jamais entre nœuds
Inférence débit-critique réplication (une instance par GPU) si le modèle tient

Le cas souvent oublié

Si le modèle tient sur un seul GPU, la meilleure stratégie de service en débit est souvent la réplication : \(N\) instances indépendantes sur \(N\) GPU, sans aucune communication.

Le parallélisme de tenseurs n'est utile que si le modèle ne tient pas, ou si la latence par requête compte plus que le débit total.

Résumé du chapitre

À retenir

  • Cinq parallélismes : données, tenseurs, pipeline, experts, séquence. On les combine.
  • Le TP enchaîne colonnes puis lignes pour n'avoir que deux all-reduce par couche. Il ne passe pas l'échelle au-delà d'un nœud (facteur 18 entre NVLink et InfiniBand).
  • NCCL détecte la topologie et choisit anneau ou arbre. Identité utile : AllReduce = ReduceScatter + AllGather.
  • NVSHMEM permet la communication initiée depuis le noyau, donc le recouvrement fin. C'est la brique de DeepEP et des megakernels multi-GPU.
  • Le contrôle au niveau du noyau permet d'inventer des motifs absents de NCCL — la « transposition distribuée » de Hazy Research divise le trafic réseau par 8.
  • Hiérarchie : HBM (3,3-8 To/s) ≫ NVLink (0,9-1,8 To/s) ≫ PCIe/IB (50-64 Go/s).
  • Si le modèle tient sur un GPU, répliquer bat souvent le parallélisme.

Vérifiez que vous avez compris

Pourquoi enchaîner « colonnes puis lignes » divise-t-il par deux le nombre d'all-reduce ?

Parce que la sortie partitionnée du découpage en colonnes est exactement l'entrée partitionnée qu'attend le découpage en lignes.

Dans un MLP \(\mathbf{y} = \sigma(\mathbf{x}\mathbf{W}_1)\mathbf{W}_2\) :

  • \(\mathbf{W}_1\) découpée en colonnes → chaque GPU a une part de \(\sigma(\mathbf{x}\mathbf{W}_1)\), sans communication ;
  • \(\mathbf{W}_2\) découpée en lignes → chaque GPU calcule une somme partielle à partir de sa part, sans avoir besoin des autres parts ;
  • un seul all-reduce à la fin.

Si l'on découpait les deux en colonnes, il faudrait un all-gather entre les deux plus un all-reduce à la fin. Le choix colonnes-puis-lignes fait disparaître la communication intermédiaire.

Vous servez un modèle de 13 milliards de paramètres sur 8 H100. TP8 ou 8 répliques ?

Cela dépend de ce que vous optimisez.

8 répliques (le modèle de 26 Go tient largement sur 80 Go) :

  • débit total maximal, aucune communication ;
  • latence par requête = celle d'un seul GPU ;
  • 8 caches KV indépendants, donc capacité de lot totale plus élevée.

TP8 :

  • latence par requête divisée par ~6 (pas 8, à cause des all-reduce) ;
  • débit total inférieur, à cause du surcoût de communication ;
  • un seul cache KV, donc moins de requêtes concurrentes.

Verdict : répliquer pour le débit, TP pour la latence. La plupart des services choisissent la réplication ; les usages interactifs à faible latence choisissent le TP.

Pourquoi la « transposition distribuée » n'est-elle pas exprimable avec NCCL ?

Parce que NCCL fournit un catalogue fixe de motifs : all-reduce, all-gather, reduce-scatter, all-to-all, broadcast.

La transposition distribuée est une repartition : passer d'une distribution où chaque GPU détient une tranche de la dimension des têtes d'attention (tensor-parallèle) à une distribution où chaque GPU détient une tranche de la dimension du lot (data-parallèle).

C'est structurellement un all-to-all avec une permutation d'indices spécifique. On pourrait l'exprimer comme un AllToAll précédé et suivi de permutations locales, mais cela impliquerait deux passages supplémentaires en mémoire et un tampon intermédiaire complet.

Dans un megakernel, où chaque thread storer écrit directement à l'adresse distante finale au moment où il produit sa donnée, la permutation est gratuite : elle est absorbée dans le calcul d'adresse.

Chapitre suivant : 8 · Entraînement contre inférence

Sources de ce chapitre