Partie 7 · GPU pour l'intelligence artificielle¶
Le cœur applicatif de ce document. Tout ce qui précède converge ici, et la partie 8 en est la conclusion logique.
Pourquoi cette partie est la plus longue¶
Parce que l'IA est aujourd'hui le domaine où se fait l'innovation en
programmation GPU. TMA, wgmma, tcgen05, FP8, FP4, CUDA Tile, Triton,
megakernels : toutes ces technologies ont été conçues pour l'apprentissage
profond, et sont ensuite adaptées ailleurs.
Comprendre l'IA du point de vue du GPU, c'est comprendre pourquoi le matériel a la forme qu'il a.
Le fil conducteur¶
Une seule question structure cette partie :
Combien d'octets faut-il déplacer pour produire un jeton ?
Toutes les techniques présentées sont des réponses à cette question :
| Technique | Ce qu'elle réduit |
|---|---|
| Pavage en tuiles (GEMM) | les relectures depuis la HBM |
| FlashAttention | la matérialisation de la matrice d'attention |
| Fusion de noyaux | les allers-retours entre opérations |
| Quantification | la taille de chaque poids |
| MoE | le nombre de poids lus |
| Recouvrement calcul/communication | le temps où rien ne circule |
| Megakernels | les bulles entre les noyaux |
Ordre de lecture¶
Séquentiel. Le chapitre 1 pose le décor, le chapitre 8 le conclut.
1 · Anatomie d'un transformeur sur GPU¶
Où passe le temps, opération par opération. Le décompte des FLOP et des octets d'une passe avant, et pourquoi ce décompte change tout entre entraînement et inférence.
2 · La GEMM, de zéro à cuBLAS¶
Les huit étapes de l'optimisation d'une multiplication matricielle, de 1 % à 95 % du pic. L'exercice canonique du domaine.
3 · FlashAttention¶
L'algorithme qui a rendu le contexte long possible. Softmax en ligne, pavage, recalcul, et l'évolution jusqu'à FlashAttention-4 sur Blackwell.
4 · La fusion de noyaux¶
Quand fusionner, quoi fusionner, et comment. Les cas rentables et ceux qui ne le sont pas.
5 · La quantification¶
FP8, FP4, INT4, formats à échelle par blocs. Ce qui se passe dans le noyau, ce que ça casse, et les kernels de référence.
6 · MoE et GEMM groupée¶
Le routage, la GEMM groupée, la répartition des experts, et les kernels de dispatch multi-GPU.
7 · Multi-GPU et collectives¶
Parallélismes de données, de tenseurs, de pipeline, d'experts. NCCL, NVSHMEM, et le recouvrement calcul/communication.
8 · Entraînement contre inférence¶
Deux régimes structurellement différents. Le décodage à lot 1 comme problème de bande passante, et pourquoi cela mène directement aux megakernels.
Le tableau qu'il faut avoir en tête¶
Pour un transformeur de \(\Theta\) paramètres, contexte \(S\), dimension cachée \(d\), lot \(b\) :
| Régime | FLOP | Octets | \(I\) | Limité par |
|---|---|---|---|---|
| Entraînement | \(\approx 6 \Theta \cdot b S\) | \(\approx 2\Theta + \text{activations}\) | élevée | le calcul |
| Préremplissage (prefill) | \(\approx 2 \Theta \cdot b S\) | \(\approx 2\Theta\) | élevée | le calcul |
| Décodage, lot 1 | \(\approx 2\Theta\) | \(\approx 2\Theta\) | ≈ 1 | la mémoire |
| Décodage, lot \(b\) | \(\approx 2\Theta b\) | \(\approx 2\Theta\) | \(\approx b\) | dépend de \(b\) |
La ligne qui explique la partie 8
Décodage à lot 1 : \(I \approx 1\), contre un seuil machine de ~296 sur H100 en BF16.
On est 296 fois du mauvais côté du roofline. Le GPU est réduit au rôle de lecteur de mémoire. Dans ce régime, chaque microseconde de surcoût — chaque lancement de noyau, chaque bulle de pipeline — est directement du temps perdu.
C'est tout le sujet des megakernels.
Sur les chiffres de cette partie
Les chiffres de performance cités proviennent des sources indiquées : papiers, documentations, blogs des équipes concernées. Ils n'ont pas été reproduits dans l'environnement de rédaction.
Les décomptes de FLOP et d'octets, en revanche, sont des calculs élémentaires que vous pouvez refaire — et que vous devriez refaire pour votre modèle.