Aller au contenu

5 · Event Tensor et le dynamisme

La limite commune à tous les megakernels précédents : ils sont statiques. Ce chapitre explique pourquoi c'est un problème sérieux, et comment l'abstraction Event Tensor l'attaque.

5.1 Le problème

Les megakernels des chapitres 3 et 4 pré-planifient tout : la liste des instructions, ou le ttGraph, est construit avant le lancement.

Or les charges réelles sont dynamiques.

Source de dynamisme Exemple
Taille de lot variable continuous batching : les requêtes arrivent et partent en permanence
Longueur de contexte variable chaque requête a son propre historique
Routage MoE le nombre de jetons par expert dépend des données
Décodage spéculatif le nombre de jetons acceptés varie à chaque pas
Attention creuse les blocs à calculer dépendent des scores
Sortie anticipée certains jetons terminent avant les autres

MPK reconnaît explicitement cette limite : « le compilateur effectue une optimisation par taille de lot ; plusieurs ttGraphs sont requis pour des tailles de lot dynamiques » et « il faut pré-générer des ttGraphs pour des tailles de lot représentatives plutôt qu'une compilation pleinement dynamique ».

Les conséquences pratiques

1. Explosion combinatoire. Avec 8 tailles de lot × 4 longueurs de contexte, c'est 32 megakernels à compiler et à garder en mémoire.

2. Temps de démarrage. Compiler 84 000 lignes de CUDA n'est pas instantané. Un service qui redémarre paie cette compilation.

3. Perte d'efficacité. Choisir le megakernel « le plus proche » d'une configuration réelle signifie souvent traiter du remplissage inutile.

4. Cas impossibles. Le décodage spéculatif et le routage MoE produisent un dynamisme intra-passe qu'aucune pré-compilation ne peut couvrir.

5.2 Le papier

Titre : Event Tensor: A Unified Abstraction for Compiling Dynamic Megakernel

Date : soumis le 14 avril 2026, révisé le 21 avril 2026 (arXiv:2604.13327)

Auteurs : Hongyi Jin, Bohan Hou, Guanjie Wang, Ruihang Lai, Jinqi Chen, Zihao Ye, Yaxing Cai, Yixin Dong, Xinhao Cheng, Zhihao Zhang, Yilong Zhao, Yingyi Huang, Lijie Yang, Jinchen Jiang, Gabriele Oliaro, Jianan Ji, Xupeng Miao, Vinod Grover, Todd C. Mowry, Zhihao Jia, Tianqi Chen

Plusieurs auteurs sont communs avec MPK — c'est explicitement une continuation.

Le résumé

Les charges GPU modernes, en particulier l'inférence de grands modèles de langage, souffrent des surcoûts de lancement de noyaux et d'une synchronisation grossière qui limitent le parallélisme inter-noyaux. Les techniques récentes de megakernel fusionnent plusieurs opérateurs en un seul noyau persistant pour éliminer les intervalles de lancement et exposer le parallélisme inter-noyaux, mais peinent à gérer les formes dynamiques et le calcul dépendant des données dans les charges réelles.

5.3 L'abstraction

La contribution centrale est une abstraction de compilation qui encode les dépendances entre tâches tuilées tout en supportant le dynamisme de forme et le dynamisme dépendant des données.

L'intuition qu'on peut en donner : au lieu de fixer le graphe de tâches à la compilation, on décrit les dépendances sous une forme paramétrée — un « tenseur d'événements » dont les indices dépendent de valeurs connues seulement à l'exécution.

MPK (statique)                    Event Tensor (dynamique)
──────────────                    ────────────────────────
ttGraph fixé à la compilation     structure de dépendances paramétrée
    │                                  │
    │ une compilation par              │ une compilation
    │ taille de lot                    │
    ▼                                  ▼
32 megakernels                    1 megakernel + résolution à l'exécution

Le compilateur associé, ETC (Event Tensor Compiler), applique des transformations d'ordonnancement statiques et dynamiques aux noyaux persistants.

5.4 Les résultats annoncés

Deux affirmations :

  1. une latence de service de LLM à l'état de l'art ;
  2. une réduction substantielle du surcoût de préchauffage (warmup) du système.

Le second point est celui qui distingue le travail. Le préchauffage inclut la compilation, la mise en cache des noyaux et l'échauffement du système ; sur un service qui doit démarrer vite ou s'adapter à une charge changeante, c'est un coût opérationnel réel, indépendant de la performance en régime établi.

Ce que le résumé ne chiffre pas

À la date de rédaction, le résumé public ne donne pas de facteurs d'accélération numériques précis. Les affirmations sont qualitatives (« state-of-the-art », « substantially reducing »).

C'est une limite de ce chapitre : il présente une direction de recherche dont les chiffres exacts demandent la lecture du papier complet. Les affirmations rapportées ici sont celles du résumé, pas des mesures reproduites.

5.5 Pourquoi le dynamisme est difficile, techniquement

Trois obstacles concrets, utiles à comprendre même sans le détail du papier.

Obstacle 1 : les compteurs de dépendance ont des cibles variables

Dans un megakernel statique, une tâche attend que son compteur atteigne une valeur connue :

while (compteurs[e] < VALEUR_CIBLE) { }

En dynamique, VALEUR_CIBLE dépend du nombre de tâches productrices — lui-même dépendant du routage MoE ou de la longueur de contexte. Il faut donc calculer les cibles à l'exécution, sur le GPU, avant de pouvoir attendre.

Obstacle 2 : l'allocation de mémoire partagée devient dynamique

Une tâche dont la taille de tuile dépend de la longueur de séquence ne peut pas demander un nombre de pages fixe. L'allocateur doit gérer des demandes variables, avec un risque d'interblocage si deux tâches attendent mutuellement des pages.

Obstacle 3 : l'équilibrage devient un problème d'ordonnancement en ligne

Avec des tâches de tailles imprévisibles, une distribution statique est nécessairement déséquilibrée. Il faut un ordonnancement dynamique — donc des atomiques, donc de la contention.

Le matériel commence à aider

Le Cluster Launch Control de Blackwell (voir CUDA moderne 3) fournit précisément une file de travail matérielle : un cluster demande le prochain identifiant de travail au lieu de le déduire de son blockIdx.

C'est l'obstacle 3 résolu en silicium. On peut raisonnablement anticiper que les générations suivantes en fourniront davantage — la trajectoire matérielle suit visiblement les besoins logiciels.

5.6 Les approches alternatives au dynamisme

Event Tensor n'est pas la seule réponse. Voici le paysage.

Approche Principe Compromis
Pré-compilation multiple (MPK) un megakernel par configuration simple, explosion combinatoire
Remplissage (padding) arrondir à la configuration supérieure simple, travail gaspillé
Recherche hors ligne (Ada-MK) DAG optimal déterminé à la compilation, pas de branchement à l'exécution rapide, ne gère pas le dynamisme intra-passe
Event Tensor dépendances paramétrées, résolues à l'exécution général, complexité de compilation
Hybride megakernel pour la partie statique, noyaux classiques pour la partie dynamique pragmatique, gain partiel

L'approche hybride mérite d'être soulignée : c'est celle d'Ada-MK en production, qui utilise TensorRT-LLM pour le préremplissage (dynamique en longueur) et le megakernel pour le décodage (plus régulier).

5.7 Ce que le dynamisme change pour l'avenir

L'enjeu réel

Si les megakernels restent statiques, ils resteront un composant spécialisé : rentables sur du décodage à configuration fixe, insérés comme greffon dans un moteur classique.

S'ils deviennent dynamiques, ils peuvent devenir le mode d'exécution par défaut d'un moteur d'inférence, remplaçant l'architecture « une bibliothèque de noyaux + un ordonnanceur CPU ».

C'est l'enjeu de cette ligne de recherche, et c'est pourquoi Event Tensor est important même sans chiffres spectaculaires : il attaque le verrou qui décide de l'avenir de la technique.

Deux indices vont dans ce sens.

Le matériel converge : Cluster Launch Control, et plus généralement l'évolution vers des mécanismes de distribution de travail en silicium.

Les compilateurs convergent : CUDA Tile, Triton, TVM/TIRx, MLIR — toute l'industrie construit des représentations intermédiaires par tuiles, qui sont exactement le bon niveau pour raisonner sur des tâches et leurs dépendances.

Résumé du chapitre

À retenir

  • Les megakernels des chapitres 3 et 4 sont statiques : tout est pré-planifié.
  • Les charges réelles sont dynamiques : lot variable, contexte variable, routage MoE, décodage spéculatif, attention creuse.
  • MPK y répond par la pré-compilation multiple, avec explosion combinatoire et surcoût de démarrage.
  • Event Tensor (avril 2026) propose une abstraction encodant les dépendances entre tâches tuilées tout en supportant le dynamisme de forme et de données. Son compilateur ETC applique des transformations statiques et dynamiques.
  • Résultats annoncés : latence à l'état de l'art et réduction substantielle du préchauffage — sans chiffres précis dans le résumé public.
  • Trois obstacles techniques : cibles de compteurs variables, allocation de mémoire partagée dynamique, ordonnancement en ligne.
  • Le Cluster Launch Control de Blackwell résout partiellement le troisième en matériel.
  • L'enjeu : megakernel comme composant spécialisé (si statique) ou comme mode d'exécution par défaut (si dynamique).

Vérifiez que vous avez compris

Pourquoi le décodage spéculatif est-il particulièrement hostile à un megakernel statique ?

Parce que son dynamisme est intra-passe et non prédictible.

À chaque pas, le modèle brouillon propose \(k\) jetons ; le grand modèle les vérifie ; on accepte un préfixe de longueur variable entre 1 et \(k\). Le nombre de jetons à traiter au pas suivant dépend donc du contenu de la vérification.

Un megakernel pré-compilé pour « \(k\) jetons » traiterait du remplissage la plupart du temps. Un megakernel pré-compilé pour chaque valeur de 1 à \(k\) multiplierait les compilations — et il faudrait choisir lequel lancer, ce qui exige une décision côté CPU, donc une synchronisation, donc le retour du surcoût qu'on cherchait à éviter.

La seule réponse propre est de rendre le megakernel capable de traiter un nombre variable de jetons sans recompilation.

En quoi le Cluster Launch Control ressemble-t-il au runtime de MPK ?

Les deux distribuent dynamiquement du travail à des unités persistantes.

MPK le fait en logiciel : 16 SM ordonnanceurs surveillent les événements et poussent les tâches prêtes dans les files des workers, avec un mélange JIT/AOT.

CLC le fait en matériel : un cluster demande le prochain identifiant de travail au lieu de le déduire de son blockIdx, sans coût d'atomiques ni de SM dédiés.

C'est une convergence significative : elle suggère que les concepteurs de puces observent ce que fait le logiciel et le câblent. Si la tendance se confirme, écrire un megakernel devrait devenir plus simple au fil des générations, pas plus difficile.

Une approche hybride (megakernel pour le décodage, moteur classique pour le préremplissage) est-elle un aveu d'échec ?

Non — c'est probablement la bonne architecture.

Le raisonnement du chapitre 1 le montre : le gain d'un megakernel est proportionnel à la part du surcoût dans le temps total. Cette part est de ~40 % en décodage à lot 1 et de moins de 1 % en préremplissage.

Mega-noyauter le préremplissage coûterait un effort d'ingénierie important pour un gain nul, tout en renonçant aux GEMM de pointe de cuBLAS et TensorRT-LLM.

L'architecture hybride applique donc chaque technique là où elle est rentable. C'est ce que fait Ada-MK en production, et c'est un signe de maturité plutôt que de renoncement.

Chapitre suivant : 6 · Fleet et les chiplets

Sources de ce chapitre