3 · L'interpréteur sur GPU¶
L'approche de Hazy Research (Stanford) : construire, à l'intérieur du GPU, un interpréteur qui exécute une séquence d'instructions. C'est le premier megakernel de modèle complet publié, et le plus pédagogique.
3.1 Le point de départ¶
Le billet Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B, publié le 27 mai 2025, part d'un constat simple.
Une passe avant de Llama-1B se décompose en une centaine de noyaux. Chacun apporte trois coûts, détaillés au chapitre 1 :
- l'ordonnancement séquentiel des lancements ;
- le surcoût de lancement (~1,3 µs même avec CUDA Graphs) ;
- les latences de chargement en début de noyau.
L'objectif annoncé : « faire tourner la passe avant d'un modèle de langage de plus d'un milliard de paramètres en 16 bits en moins d'une milliseconde sur un GPU » — ce qui, selon les auteurs, n'avait jamais été fait.
3.2 Le problème 1 : fusionner des dizaines d'opérations¶
La solution : un interpréteur¶
Chaque SM reçoit une séquence d'instructions pré-planifiée. Un interpréteur embarqué les exécute l'une après l'autre.
Sept instructions fusionnées suffisent pour la passe avant de Llama :
| # | Instruction |
|---|---|
| 1 | RMSNorm + QKV + RoPE fusionnés |
| 2 | Calcul de l'attention |
| 3 | Réduction de l'attention |
| 4 | Projection O + résidu |
| 5 | RMSNorm + up/gate + SiLU fusionnés |
| 6 | Projection down + résidu |
| 7 | RMSNorm + tête de modélisation du langage |
Chaque instruction est écrite selon un gabarit CUDA commun, ce qui garantit qu'elles interopèrent dans l'interpréteur.
// Structure conceptuelle du gabarit
template <typename Config>
struct Instruction {
static __device__ void loader(state& s, const args& a); // TMA vers les pages
static __device__ void compute(state& s, const args& a); // MMA, arithmétique
static __device__ void storer(state& s, const args& a); // écriture, compteurs
};
C'est la structure loader / compute / storer de la
spécialisation des warps,
appliquée à chaque instruction.
3.3 Le problème 2 : les bulles de pipeline mémoire¶
Le constat¶
On veut charger les poids de l'instruction \(n+1\) pendant que l'instruction \(n\) termine. Mais la mémoire partagée est occupée par l'instruction \(n\).
La solution : un allocateur de pages¶
Les 213 Ko de mémoire partagée d'un H100 sont découpés en treize pages de 16 Ko.
┌──────┬──────┬──────┬──────┬──────┬──────┬──────┬─────┐
│ p0 │ p1 │ p2 │ p3 │ p4 │ ... │ p11 │ p12 │
│16 Ko │16 Ko │16 Ko │16 Ko │16 Ko │ │16 Ko │16Ko │
└──────┴──────┴──────┴──────┴──────┴──────┴──────┴─────┘
↑ ↑ ↑ ↑
instr.n instr.n instr.n instr. n+1
(déjà en chargement)
Chaque instruction demande explicitement des pages et les libère dès qu'elle n'en a plus besoin. L'interpréteur passe automatiquement les pages libérées à l'instruction suivante.
Le résultat : dès qu'une page se libère, le chargement des poids suivants commence, sans attendre la fin de l'instruction courante.
Pourquoi cette idée est importante
Elle transforme la mémoire partagée d'une variable statique en une ressource dynamique gérée. C'est le seul moyen de faire du pipelining entre opérations de tailles et de natures différentes.
Mirage MPK a adopté la même idée, avec des pages de 32 Ko et une « abstraction de mémoire partagée paginée permettant le partage de ressources entre tâches ».
3.4 Le problème 3 : la synchronisation sans frontières de noyau¶
La solution : des compteurs en mémoire globale¶
Avant le lancement, un tableau de compteurs est initialisé à zéro. Chaque instruction incrémente son compteur en terminant ; les instructions dépendantes attendent que leurs compteurs atteignent une valeur cible.
L'exemple canonique du découpage¶
C'est le passage le plus instructif du billet.
Le MLP produit un état caché, consommé par la projection descendante. Avec PDL, il faudrait attendre tout l'état caché.
Leur solution : produire et consommer l'état caché en quatre morceaux, avec un compteur par morceau.
Sans découpage (PDL) :
up/gate ████████████████
down ████████████████
Avec 4 compteurs :
up/gate ████│████│████│████
down ████│████│████│████
↑ démarre dès que le 1er quart est prêt
La profondeur du chemin critique est réduite de ~40 % sur cette portion.
La généralisation
Cet exemple illustre le principe général : la granularité de la dépendance doit correspondre à la granularité de la production, pas à la granularité de l'opération.
C'est précisément ce qu'un compilateur peut automatiser — d'où Mirage MPK, dont la « décomposition d'opérateurs » partitionne les sorties entre SM et dont l'« analyse de dépendances » énumère les paires de tâches et introduit des événements pour les régions de données qui se recouvrent.
3.5 Les résultats du megakernel de latence¶
| Plateforme | Temps de passe avant | Contre vLLM | Contre SGLang |
|---|---|---|---|
| H100 | ~1 ms | 2,5× | 1,5× |
| B200 | ~680 µs | 3,5× | 1,5× |
Utilisation de la bande passante mémoire : 78 %, contre ~50 % pour les systèmes existants.
Les auteurs indiquent qu'il s'agit de « la première fois que la passe avant d'un modèle de langage de plus d'un milliard de paramètres en 16 bits est exécutée en moins d'une milliseconde sur un GPU » (sur H100).
3.6 Le megakernel de débit¶
Quatre mois plus tard, en septembre 2025, l'équipe publie We Bought the Whole GPU, So We're Damn Well Going to Use the Whole GPU : un megakernel orienté débit, pour Llama-70B en tensor-parallèle sur 8 H100.
Pourquoi c'est un objet différent¶
| Megakernel latence | Megakernel débit | |
|---|---|---|
| Modèle | Llama-1B | Llama-70B |
| GPU | 1 | 8, tensor-parallèles |
| Lot | 1 | jusqu'à 8 192 |
| Objectif | supprimer les frontières | saturer des ressources hétérogènes |
| Goulot | latence mémoire | équilibre entre tensor cores, mémoire, NVLink |
Le point clé : à grand lot, les frontières de noyau sont amorties. Ce qui compte alors est que les différentes ressources matérielles soient utilisées simultanément — les GEMM sont limitées par le calcul, le décodage d'attention par la mémoire, la communication par NVLink.
Le recouvrement à trois niveaux¶
1. À l'intérieur d'un SM. Threads loader, compute et storer spécialisés, permettant de charger les poids de l'opération suivante pendant qu'on termine la courante.
→ 2 à 6 % de gain.
2. Entre les SM. Une file de travail globale distribue dynamiquement les instructions, et l'entrelacement d'instructions d'opérations différentes permet de faire tourner simultanément des tâches à dominante calcul et à dominante communication sur des SM différents.
→ 14,2 % de gain à lot 8 192, contre un ordonnancement en tourniquet.
3. Entre les GPU. Des threads storer dédiés gèrent la communication inter-GPU de manière asynchrone via une abstraction de « Parallel Global Layout », libérant les autres threads.
La transposition distribuée¶
L'optimisation la plus originale, déjà évoquée en IA 7.
Pour réduire le trafic réseau d'un facteur huit, les auteurs répliquent la
matrice de projection O sur tous les GPU et l'exécutent en parallélisme de
données, ce qui élimine le reduce-scatter post-attention. Il est remplacé
par une « transposition distribuée » qui repartitionne les données du format
tensor-parallèle vers le format data-parallèle.
Ils notent que cette opération est « facilement exprimable dans le cadre du megakernel » mais « pas efficacement exprimable avec les motifs de communication standard ».
Les neuf instructions de Llama-70B¶
- RMSNorm + all-gather
- Projection QKV + RoPE
- Attention + transposition distribuée
- Projection O
- Gate + SiLU
- Projection up
- Projection down + reduce-scatter
- RMSNorm finale
- Tête de modélisation du langage
Les résultats¶
Sur 65 536 invites de ShareGPT, 8 H100 :
| Métrique | Megakernel | SGLang |
|---|---|---|
| Débit d'entrée | 14 425 jetons/s | 11 783 jetons/s |
| Débit de sortie | 9 043 jetons/s | 7 387 jetons/s |
| Débit total | 23 468 jetons/s | 19 170 jetons/s |
Soit +22 %.
Le megakernel est intégré à Tokasaurus, un moteur d'inférence qui gère l'ordonnancement des lots et les pages du cache KV. L'ordonnanceur CPU génère les instructions sur 64 threads avec plus de 90 % de temps d'inactivité CPU, ce qui montre que la coordination hôte-périphérique n'est pas limitante.
Les auteurs reconnaissent qu'il reste des optimisations possibles, citant notamment les « débordements de registres et autres problèmes bas niveau », et présentent les 22 % comme un point de départ plutôt qu'un plafond.
3.7 Ce que ThunderKittens apporte¶
Les deux megakernels sont écrits avec ThunderKittens, la bibliothèque de tuiles de la même équipe.
Sans elle, chaque instruction demanderait d'écrire à la main les descripteurs
TMA, les dispositions wgmma et le swizzling. Avec elle, une instruction
ressemble à l'algorithme :
rt_fl<16, 64> att_block;
mma_ABt(att_block, q_reg, k_smem, att_block);
row_max(max_vec, att_block, max_vec);
exp(att_block, att_block);
C'est ce qui rend un megakernel de plusieurs milliers de lignes relisible.
3.8 Ce que cette approche montre et ce qu'elle ne montre pas¶
Ce qu'elle démontre
- Un megakernel de modèle complet est réalisable et apporte un gain substantiel (2,5× sur vLLM, 1,5× sur SGLang à lot 1).
- Le plafond de bande passante peut être poussé de ~50 % à 78 %.
- L'approche s'étend au débit et au multi-GPU, avec des gains différents mais réels (+22 %).
- L'allocateur de pages et les compteurs fins sont les deux mécanismes décisifs.
Ce qu'elle ne démontre pas
- La généralité. Les instructions sont écrites à la main pour Llama. Changer d'architecture demande de réécrire.
- La maintenabilité. Un megakernel écrit à la main est un objet de recherche, difficile à faire évoluer.
- Le dynamisme. Les instructions sont pré-planifiées ; formes variables et dépendances aux données ne sont pas traitées.
Ces trois limites sont précisément ce qu'attaquent Mirage MPK (généralité par compilation) et Event Tensor (dynamisme).
Résumé du chapitre¶
À retenir
- Hazy Research construit un interpréteur sur GPU : chaque SM exécute une
séquence d'instructions pré-planifiées, chacune suivant un gabarit
loader/compute/storer. - Sept instructions suffisent pour Llama-1B, neuf pour Llama-70B en tensor-parallèle.
- L'allocateur de pages (13 pages de 16 Ko sur H100) transforme la mémoire partagée en ressource dynamique et permet le pipelining entre instructions.
- Les compteurs en mémoire globale remplacent les frontières de noyau, avec un découpage en quatre morceaux là où PDL forcerait à tout attendre.
- Résultats latence : < 1 ms sur H100, ~680 µs sur B200, 78 % de la bande passante, 2,5× vLLM et 1,5× SGLang.
- Résultats débit : +22 % sur SGLang (23 468 contre 19 170 jetons/s) sur Llama-70B/8×H100, avec un recouvrement à trois niveaux et une « transposition distribuée » qui divise le trafic réseau par 8.
- Limites : écrit à la main, spécifique à Llama, statique.
Vérifiez que vous avez compris¶
Pourquoi 13 pages de 16 Ko plutôt qu'une allocation contiguë par instruction ?
Parce qu'une allocation contiguë interdit le pipelining.
Si l'instruction \(n\) occupe les 213 Ko et ne les libère qu'à sa fin, l'instruction \(n+1\) ne peut rien charger avant. La bulle mémoire réapparaît, exactement comme entre deux noyaux.
Avec des pages, l'instruction \(n\) libère ses pages au fur et à mesure — dès qu'elle a fini de lire une tuile de poids, la page est rendue —, et l'interpréteur y lance immédiatement le chargement de l'instruction suivante.
La granularité de 16 Ko est un compromis : assez petite pour libérer tôt, assez grande pour qu'une copie TMA soit efficace.
Le megakernel de débit gagne 14,2 % par l'ordonnancement inter-SM, contre 2-6 % par le pipelining intra-SM. Pourquoi cet écart ?
Parce que les deux corrigent des problèmes d'échelles différentes.
Le pipelining intra-SM récupère des bulles courtes — quelques dizaines de cycles pendant lesquels un SM attend une donnée. Avec un pipeline déjà profond, ces bulles sont rares.
L'ordonnancement inter-SM corrige le déséquilibre de charge. À lot 8 192, avec des opérations très hétérogènes (GEMM limitée par le calcul, attention limitée par la mémoire, communication limitée par NVLink), un ordonnancement statique laisse des SM inactifs pendant de longues périodes.
C'est la même leçon que dans tout ce document : les déséquilibres et les frontières coûtent plus que les micro-inefficacités.
Le CPU génère les instructions avec plus de 90 % de temps d'inactivité. Pourquoi est-ce important ?
Parce que cela répond à une objection naturelle : « si le GPU exécute un interpréteur, le CPU ne devient-il pas le goulot en générant le programme ? »
Non. La génération d'instructions est un travail léger — construire des structures de quelques centaines d'octets décrivant des tâches — et elle est pipelinée avec l'exécution GPU.
Ce chiffre valide aussi le choix architectural : puisque le CPU a de la marge, on peut lui confier davantage (ordonnancement adaptatif, gestion de lots dynamiques) sans risquer de créer un goulot.
À titre de comparaison, MPK stocke chaque description de tâche sur 352 octets en mémoire du périphérique, ce qui donne l'ordre de grandeur du volume à produire.
Chapitre suivant : 4 · Mirage Persistent Kernel