9 · Quand ne pas le faire¶

Le chapitre qui manque à la plupart des articles sur le sujet. Un megakernel est une technique puissante et étroite : hors de son domaine, c'est du travail perdu et de la dette technique.

9.1 Le calcul de rentabilité, formalisé¶

Soit \(t_{\min}\) le plancher physique (octets à lire ÷ bande passante) et \(t_{\text{réel}}\) le temps mesuré.

\[ \text{surcoût} = \frac{t_{\text{réel}} - t_{\min}}{t_{\text{réel}}} \]

Un megakernel ne peut récupérer qu'une fraction de ce surcoût. Les données publiées donnent un ordre de grandeur :

Hazy Research : de ~50 % à 78 % de la bande passante, soit ~56 % du surcoût récupéré ;
MPK sur Qwen3-8B/A100 : 14,5 → 12,5 ms avec un plancher à ~10 ms, soit 44 % du surcoût récupéré.

Estimation prudente : un megakernel récupère 40 à 55 % du surcoût.

D'où :

\[ \text{gain attendu} \approx 0{,}5 \times \text{surcoût} \]

Situation	Surcoût	Gain attendu
Entraînement, gros lot	< 1 %	négligeable
Préremplissage	~2 %	négligeable
Décodage, 70 B, lot 64	~10 %	~5 %
Décodage, 8 B, lot 8	~25 %	~12 %
Décodage, 1 B, lot 1	~45 %	~25 %

Le seuil de décision

En dessous de 20 % de surcoût, le gain attendu est inférieur à 10 %.

Compte tenu du coût d'ingénierie — plusieurs semaines à plusieurs mois, plus une dette de maintenance permanente — ce n'est pas rentable.

9.2 Les six situations où il ne faut pas¶

1. Le régime est limité par le calcul¶

Entraînement, préremplissage, inférence à grand lot : \(I \gg I_{\text{crit}}\), les tensor cores sont saturés, et le surcoût de lancement est noyé.

Un megakernel n'y apporte rien, et il vous prive de cuBLAS et cuDNN, qui sont meilleurs que ce que vous écrirez.

2. Le modèle change souvent¶

Un megakernel encode l'architecture du modèle dans son jeu d'instructions. Changer de modèle demande de réécrire ou de recompiler.

Si vous itérez sur l'architecture — recherche, expérimentation, prototypage — c'est le pire moment pour figer le chemin d'exécution.

3. Les formes sont très dynamiques¶

C'est la limite du chapitre 5. MPK reconnaît devoir « pré-générer des ttGraphs pour des tailles de lot représentatives ».

Avec 8 tailles de lot × 4 longueurs de contexte, vous compilez 32 megakernels. Le temps de compilation et la mémoire occupée deviennent le problème.

4. L'équipe n'a pas l'expertise¶

Il faut maîtriser : le modèle mémoire de CUDA, la synchronisation inter-blocs, les copies asynchrones, TMA, wgmma, la spécialisation des warps, et le débogage de courses de données sur GPU.

Le facteur bus

Un megakernel écrit par une personne qui part est un actif qui devient un passif. Il est illisible pour qui n'a pas ce contexte, et il touche à tout le modèle — donc on ne peut pas le contourner.

Un billet de blog de 2025 sur le sujet est intitulé, non sans raison, How to build unmaintainable kernels.

5. Le goulot est ailleurs¶

Avant d'optimiser l'exécution du modèle, vérifiez que c'est bien elle qui limite. Les goulots réels d'un service d'inférence :

Goulot	Fréquence
Tokenisation et prétraitement	fréquent
Ordonnancement des requêtes	fréquent
Gestion du cache KV (fragmentation, évictions)	fréquent
Sérialisation réseau	fréquent
Exécution du modèle	pas toujours dominant

Un profilage de bout en bout est un prérequis.

6. Une solution plus simple existe¶

Par ordre croissant d'effort :

torch.compile(mode="reduce-overhead") — CUDA Graphs, gratuit ;
quantification — divise le trafic par 2 ou 4, c'est le levier direct ;
fusion ciblée des opérations élémentaires — Liger Kernel, Triton ;
décodage spéculatif — augmente le travail utile par passe ;
regroupement — si la latence le permet ;
MoE — si vous contrôlez l'architecture ;
essayer MPK — un compilateur existe, utilisez-le ;
seulement ensuite : écrire un megakernel.

9.3 Les coûts, en détail¶

Le coût de développement¶

Approche	Ordre de grandeur
Utiliser MPK	quelques jours
Adapter un megakernel existant à un modèle proche	2 à 6 semaines
Écrire de zéro pour une nouvelle architecture	2 à 6 mois

Pour situer, MPK représente ~40 000 lignes de C++, ~84 000 lignes de CUDA et ~10 000 lignes de Python — plusieurs années-personnes.

Le coût de compilation¶

Un noyau CUDA unique contenant tout le modèle est long à compiler. Les templates, le déroulage et l'inlining produisent des unités de compilation énormes, ce qui allonge le cycle d'itération de manière très concrète.

Le coût en occupancy¶

Un megakernel occupe généralement un bloc par SM, avec toute la mémoire partagée réservée pour les pages. MPK utilise 128 threads par worker sur H100, soit 4 warps sur 64 — environ 6 % d'occupancy.

C'est acceptable si et seulement si le pipelining logiciel est excellent. Si ce n'est pas le cas, vous perdez à la fois le TLP et le bénéfice de la fusion.

Le coût en registres¶

Le compilateur alloue les registres pour le pire chemin du switch. Une seule instruction gourmande impose sa consommation à toutes les autres, et le spilling devient facile.

Hazy Research cite explicitement les « débordements de registres et autres problèmes bas niveau » parmi les optimisations restantes de leur megakernel de débit — dans un travail qui atteint pourtant +22 %.

Le coût de débogage¶

Une course de données entre deux instructions d'un megakernel se manifeste par un jeton incorrect une fois sur mille. C'est le type de bug le plus coûteux qui soit.

9.4 Ce que dit la littérature, honnêtement¶

Le papier fondateur de 2012 sur les persistent threads concluait déjà :

L'approche PT peut atteindre jusqu'à un ordre de grandeur d'accélération sur les noyaux non-PT, mais peut aussi entraîner une perte de performance dans de nombreux cas.

Quatorze ans plus tard, MPK annonce « 1,0 – 1,7× » contre SGLang et vLLM. La borne basse de 1,0 signifie : sur certaines configurations, aucun gain.

C'est l'information la plus utile de tout ce chapitre, et elle vient des auteurs eux-mêmes.

9.5 Les alternatives partielles¶

Si vous voulez une partie du bénéfice sans le coût complet.

Technique	Gain typique	Effort
CUDA Graphs	1,1-1,4× sur les noyaux courts	trivial
Fusion des opérations élémentaires	1,2-2×	faible
Fusion par couche (un noyau par couche)	1,1-1,3×	moyen
PDL sur les frontières critiques	1,05-1,15×	moyen
MPK (compilateur existant)	1,0-1,7×	faible
Megakernel manuel	1,5-2,5×	très élevé

La ligne à retenir

Essayez MPK. C'est le seul moyen d'obtenir un vrai megakernel pour un effort raisonnable, et ses auteurs annoncent la compilation d'un modèle Hugging Face « en quelques dizaines de lignes de Python ».

Si MPK ne couvre pas votre cas et que le gain justifie plusieurs mois de travail, alors écrivez le vôtre. Pas avant.

9.6 L'arbre de décision¶

Mon inférence est trop lente
   │
   ├─ Ai-je profilé de bout en bout ?
   │     NON → profiler. Le goulot est souvent hors du modèle.
   │
   ├─ Quel est le surcoût ? (t_réel − t_min) / t_réel
   │     < 20 % → un megakernel n'apportera pas 10 %. STOP.
   │
   ├─ Ai-je essayé, dans l'ordre :
   │     · CUDA Graphs                    → gratuit
   │     · quantification                 → facteur 2-4
   │     · fusion (Liger, torch.compile)  → 1,2-2×
   │     · décodage spéculatif            → 1,5-2,5×
   │     · regroupement                   → si la latence le permet
   │     NON → les faire d'abord.
   │
   ├─ Mon modèle est-il stable ?
   │     NON → attendre qu'il le soit.
   │
   ├─ Mes formes sont-elles stables ?
   │     NON → approche hybride, ou attendre Event Tensor.
   │
   ├─ MPK couvre-t-il mon architecture ?
   │     OUI → l'utiliser. FIN.
   │
   ├─ Ai-je 2 à 6 mois et l'expertise ?
   │     NON → STOP.
   │
   └─ OUI → écrire le megakernel. Voir le chapitre 8.

Résumé du chapitre¶

À retenir

Un megakernel récupère 40 à 55 % du surcoût, d'après les données publiées. En dessous de 20 % de surcoût, ce n'est pas rentable.
Six contre-indications : régime limité par le calcul, modèle instable, formes très dynamiques, expertise absente, goulot ailleurs, solution plus simple disponible.
Coûts : 2 à 6 mois de développement de zéro, compilation lente, occupancy à ~6 %, registres alloués pour le pire chemin, débogage de courses.
Le papier de 2012 disait déjà que les persistent threads « peuvent aussi entraîner une perte de performance ». MPK annonce une borne basse de 1,0×.
Essayez MPK avant d'écrire. Si le compilateur couvre votre cas, le rapport effort/gain est incomparable.

Vérifiez que vous avez compris¶

Votre service d'inférence sert un modèle de 70 B à lot 128. Un megakernel vaut-il le coup ?

Non, presque certainement.

À lot 128, \(I \approx 128\) contre un seuil de 296 : on est encore limité par la mémoire mais pas d'un facteur écrasant, et surtout le travail par passe est important — de l'ordre de 25 à 40 ms.

Le surcoût de lancement (560 noyaux × 1,3 µs ≈ 0,7 ms) représente ~2 à 3 %. Même en récupérant la moitié, le gain est de ~1,5 %.

Les leviers rentables à ce lot : quantification (facteur 2 sur le trafic), optimisation du cache KV, meilleur ordonnancement des requêtes, désagrégation préremplissage/décodage.

Pourquoi le coût en registres est-il structurellement plus élevé dans un megakernel ?

Parce que le compilateur alloue les registres statiquement, pour la totalité du noyau.

Dans une architecture à noyaux multiples, chaque noyau a sa propre allocation : le noyau de RMSNorm en utilise 32, celui de GEMM 168. Chacun est optimal pour lui-même.

Dans un megakernel, le switch contient les deux chemins. Le compilateur doit garantir que le chemin le plus gourmand a assez de registres, donc il alloue 168 registres à tout le noyau — y compris pendant l'exécution de RMSNorm.

Conséquences : occupancy plafonnée par le pire chemin, et risque de spilling si un chemin dépasse 255 registres.

Atténuations : __launch_bounds__, découper les instructions les plus gourmandes, ou utiliser setmaxnreg (Hopper) pour redistribuer entre warp groups.

MPK annonce « 1,0-1,7× ». Comment décider si votre cas sera près de 1,0 ou de 1,7 ?

Faites le calcul du surcoût avant d'essayer.

Comptez les octets minimaux : poids actifs + cache KV lu par pas.
Divisez par la bande passante de votre carte → \(t_{\min}\).
Mesurez \(t_{\text{réel}}\) avec votre moteur actuel.
\(\text{surcoût} = (t_{\text{réel}} - t_{\min}) / t_{\text{réel}}\).
surcoût < 15 % → attendez-vous à ~1,0-1,05× ;
surcoût ~30 % → ~1,15-1,2× ;
surcoût > 45 % → ~1,3-1,7×.

Ce calcul prend dix minutes et vous dit à quoi vous attendre. C'est exactement le raisonnement du modèle roofline, appliqué à un système entier plutôt qu'à un noyau.

Chapitre suivant : 10 · État de l'art et perspectives

Sources de ce chapitre¶

Gupta, Stuart, Owens, A Study of Persistent Threads Style GPU Programming, InPar 2012 — eScholarship
Mirage Persistent Kernel, arXiv:2512.22219 — « 1,0-1,7× », limites reconnues, 14,5 → 12,5 ms avec plancher à 10 ms.
Hazy Research, Look Ma, No Bubbles! et We Bought the Whole GPU — 78 % de bande passante, débordements de registres cités comme limite.
How to build unmaintainable kernels, Ian Barber