Résumé exécutif¶

Ce que ce document établit, en une page dense. Chaque affirmation renvoie au chapitre qui la démontre.

Le GPU en une idée¶

Un CPU est optimisé pour finir vite une tâche ; un GPU est optimisé pour en finir beaucoup. Toute l'architecture découle de ce choix : peu de cache par thread, pas de prédiction de branchement sophistiquée, mais des dizaines de milliers de threads en vol simultanément dont le matériel commute le contexte à chaque cycle. La latence n'est pas réduite — elle est cachée.

Conséquence pratique, qui surprend tout débutant : un GPU ne va pas vite parce qu'il calcule vite, il va vite parce qu'il attend efficacement. → Fondations 1, Fondations 3

Le chiffre qui décide de tout¶

Pour un noyau donné, on calcule l'intensité arithmétique \(I\) — le nombre d'opérations effectuées par octet lu ou écrit en mémoire globale :

\[ I = \frac{\text{opérations}}{\text{octets transférés}} \]

et on la compare au ratio machine \(I_{\text{crit}} = P / B\), où \(P\) est le débit de calcul crête et \(B\) la bande passante mémoire.

Si \(I < I_{\text{crit}}\) : le noyau est limité par la mémoire. Optimiser le calcul ne sert à rien.
Si \(I > I_{\text{crit}}\) : il est limité par le calcul.

Sur un B200 en BF16, \(I_{\text{crit}} \approx 2\,250 / 8 \approx 280\) opérations par octet. La quasi-totalité des noyaux d'inférence LLM à petit lot sont très en dessous. C'est pourquoi l'inférence LLM est un problème de bande passante mémoire, pas de puissance de calcul. → Performance 1, IA 8

Les six leviers d'optimisation, par ordre de rendement¶

Ne pas transférer — fusionner les opérations pour éviter les aller-retours en mémoire globale. Le levier le plus rentable, de loin.
Transférer bien — accès coalescés, pas de conflit de banc en mémoire partagée. Un facteur 8 à 32 est en jeu.
Utiliser les unités matricielles — tensor cores : facteur 8 à 30 sur les produits matriciels.
Recouvrir — charger le bloc \(n+1\) pendant qu'on calcule le bloc \(n\). C'est le pipelining logiciel.
Réduire la précision — FP16, FP8, FP4 : chaque division par deux de la taille double la bande passante effective et le débit tensor core.
Occuper — assez de warps pour cacher la latence, mais pas trop, sinon on perd des registres. L'occupancy est un moyen, jamais un objectif.

→ Performance 5

Ce qui a changé depuis 2022¶

Hopper (2022) puis Blackwell (2024-2025) ont déplacé le centre de gravité de la programmation GPU du thread vers la tuile asynchrone :

Mécanisme	Ce qu'il change	Génération
`cp.async`	copie globale → partagée sans passer par les registres	Ampere
TMA	un seul thread déclenche la copie d'un tenseur jusqu'à 5D	Hopper
Clusters	plusieurs blocs synchronisés sur des SM voisins, mémoire partagée distribuée	Hopper
`wgmma`	produit matriciel à l'échelle du warp group (128 threads)	Hopper
`tcgen05`	produit matriciel jusqu'à 256×256×16 sur une paire de SM, opérandes en tensor memory	Blackwell
TMEM	256 Ko de mémoire dédiée aux tensor cores par SM	Blackwell
CUDA Tile	modèle de programmation par tuiles, le compilateur place les threads	CUDA 13.1

Sur Blackwell, une instruction tcgen05.mma a une latence quasi constante de ≈ 11 cycles, contre 32 à 128 cycles pour wgmma selon la taille de tuile — soit 2,9 à 11,6× de moins (microbenchmarking Blackwell, arXiv:2512.02189). → CUDA moderne

Le focus : les megakernels¶

Le problème¶

Une passe avant de Llama-1B, c'est une centaine de lancements de noyaux. Chaque frontière de noyau coûte trois choses :

une barrière globale implicite — aucun bloc du noyau \(n+1\) ne démarre avant que tous les blocs du noyau \(n\) soient finis. Avec 512 blocs sur 148 SM, 80 SM attendent le retardataire ;
le lancement lui-même — ≈ 1,3 µs même avec CUDA Graphs (Hazy Research), ≈ 4,5 µs mesuré sur MI300X (Kog) ;
une bulle mémoire — au début de chaque noyau, plus rien n'est en vol ; il faut de nouveau attendre des milliers de cycles que les poids arrivent.

Sur des modèles de langage en décodage à lot 1, ces bulles dominent : la mesure de référence sur H100 montrait environ 770 passes avant par seconde réalisées contre ~1 350 permises par la bande passante.

La solution¶

Un megakernel : un unique lancement de noyau, persistant, qui exécute le modèle entier. Les dépendances ne sont plus exprimées par des frontières de noyau mais par des compteurs en mémoire globale, à une granularité beaucoup plus fine que « tout le noyau précédent ».

Trois familles d'implémentation, détaillées en partie 8 :

Approche	Idée	Résultat annoncé
Interpréteur sur GPU (Stanford, ThunderKittens)	chaque SM exécute une file d'« instructions » pré-planifiées, avec un allocateur de pages de mémoire partagée	Llama-1B : < 1 ms par passe avant sur H100, 78 % de la bande passante ; 2,5× vs vLLM, 1,5× vs SGLang
Compilateur (Mirage MPK, CMU)	graphe de tâches au niveau SM, événements, ordonnanceurs embarqués	1,0–1,7× vs SGLang/vLLM ; Qwen3-8B sur A100 : 14,5 ms → 12,5 ms par jeton
Monokernel matériel-spécifique (Kog, AMD MI300X)	synchronisation par sentinelle NaN plutôt que par atomiques	synchronisation de grille : 7,6–7,9 µs → 0,8–0,9 µs

Où en est-on en 2026¶

Le sujet est passé du prototype à la production. Ada-MK est déployé dans le système publicitaire de Baidu ; Event Tensor s'attaque au dynamisme (formes variables, dépendances aux données) ; Fleet traite le cas des GPU multi-puces ; AutoMegaKernel fait générer et vérifier statiquement les ordonnancements par un agent.

Mais le megakernel n'est pas gratuit : il faut recompiler par taille de lot, la mémoire partagée devient une ressource globale à budgéter, et le débogage d'un noyau unique de 84 000 lignes de CUDA est un problème en soi. → Megakernels 9

Comment on écrit un noyau en 2026¶

Il n'y a plus une réponse mais un gradient, du plus productif au plus contrôlé :

torch.compile  →  Helion  →  Triton  →  Gluon / TLX  →  CuTe DSL  →  CUDA C++  →  PTX  →  SASS
   automatique                                                                        illisible

Règle pratique : descendez d'un cran seulement quand le profileur vous dit pourquoi. Triton couvre bien la fusion élémentaire et l'attention ; CuTe DSL et CUTLASS couvrent la GEMM de pointe (FlashAttention-4 est écrit entièrement en CuTe DSL) ; le CUDA C++ brut reste nécessaire pour les megakernels et tout ce qui touche à la synchronisation inter-blocs. → Écosystème 1

Cinq erreurs qui coûtent le plus cher¶

Les cinq classiques

Optimiser le calcul d'un noyau limité par la mémoire. Toujours calculer \(I\) d'abord.
Viser 100 % d'occupancy. Volkov a démontré dès 2010 qu'on peut saturer une machine à 25 % d'occupancy avec assez de parallélisme d'instructions.
Mesurer sans échauffement. La première exécution inclut la compilation JIT, les défauts de page et une horloge basse.
Comparer à la mauvaise ligne de base. « 10× plus rapide que NumPy » ne dit rien ; « 0,93× de cuBLAS » dit tout.
Croire que __syncthreads() synchronise le GPU. Il synchronise un bloc. La synchronisation inter-blocs est un sujet à part entière — et c'est exactement celui des megakernels.

Ce que ce document ne tranche pas¶

Le megakernel va-t-il devenir le mode d'exécution par défaut ? Les gains sont réels à petit lot ; à grand lot ils se réduisent parce que les frontières de noyau sont amorties. La question reste ouverte.
Les LLM vont-ils écrire les noyaux ? Sur KernelBench, les modèles de pointe ne battent la ligne de base PyTorch que dans moins de 20 % des cas (KernelBench-X, arXiv:2605.04956), mais les approches agentiques avec profilage en boucle progressent vite.
La portabilité vaut-elle son coût ? L'écart mesuré entre code portable et code natif reste de 10 à 30 % au niveau du noyau. Le calcul dépend de votre parc, pas d'un principe.