Aller au contenu

2 · Définition et généalogie

Ce qu'est exactement un megakernel, ce qu'il n'est pas, et d'où vient l'idée.

2.1 La définition

Définition

Un megakernel est un noyau GPU unique et persistant qui exécute l'intégralité d'un programme tensoriel — typiquement une passe avant de modèle — en remplaçant les frontières de noyau par des dépendances à grain fin exprimées en mémoire globale.

Quatre propriétés le caractérisent.

1. Un seul lancement. Le modèle entier est exécuté par un unique appel <<<grille, bloc>>>.

2. Persistance. La grille est dimensionnée pour que tous les blocs soient simultanément résidents — typiquement un bloc par SM. Le noyau ne se termine qu'à la fin du travail.

3. Dépendances fines. Les dépendances entre opérations sont exprimées par des compteurs, des sémaphores ou des sentinelles en mémoire globale, à une granularité bien plus fine que « tout le noyau précédent ».

4. Hétérogénéité interne. Un même noyau exécute des opérations différentes — normalisations, GEMM, attention, communication — selon un mécanisme de dispatch interne.

2.2 Ce que ce n'est pas

Trois confusions fréquentes.

Ce n'est pas de la simple fusion de noyaux

La fusion combine des opérations qu'un bloc peut calculer seul. Un megakernel franchit la barrière de la synchronisation inter-blocs.

Fusion Megakernel
Portée ce qu'un bloc peut faire le modèle entier
Synchronisation __syncthreads() compteurs globaux
Dispatch aucun interpréteur
Nombre de lancements réduit un

Ce n'est pas un CUDA Graph

Un CUDA Graph reste une séquence de noyaux avec leurs barrières et leurs bulles. Il ne supprime que le surcoût CPU.

Ce n'est pas un noyau persistant ordinaire

Un noyau persistant classique (une GEMM avec ordonnanceur de tuiles, par exemple) exécute une seule opération avec une file de travail. Un megakernel exécute plusieurs opérations différentes, avec des dépendances entre elles.

La distinction est de degré autant que de nature : le noyau persistant est la brique, le megakernel est ce qu'on en fait.

2.3 La généalogie

2012 ── Persistent Threads (Gupta, Stuart, Owens, UC Davis)
   │    Le concept fondateur : garder les blocs résidents,
   │    consommer une file de travail.
   │    Quatre cas d'usage identifiés :
   │      · synchronisation CPU-GPU
   │      · équilibrage de charge / parallélisme irrégulier
   │      · localité producteur-consommateur
   │      · synchronisation globale
   │
2017 ── Cooperative Groups (CUDA 9)
   │    grid.sync() officialise la barrière inter-blocs
   │
2018 ── Persistent RNN, Persistent kernels temps réel
   │    Premiers usages en apprentissage profond
   │
2020 ── cp.async (Ampere)
   │    Copies asynchrones : le pipelining devient praticable
   │
2022 ── TMA, clusters, wgmma (Hopper)
   │    Les briques d'un pipeline profond sont réunies
   │
2023 ── FlashAttention, Stream-K, tile schedulers
   │    Les noyaux persistants deviennent standard pour UNE opération
   │
2025 ── LOOK MA, NO BUBBLES (Hazy Research, mai)
   │    Premier megakernel de modèle complet.
   │    Llama-1B en moins d'une milliseconde sur H100.
   │
2025 ── Megakernel tensor-parallèle (Hazy Research, sept.)
   │    Version orientée débit, 8 GPU, +22 % sur SGLang
   │
2025 ── Mirage Persistent Kernel (CMU et al., déc.)
   │    Premier COMPILATEUR de megakernel
   │
2026 ── Event Tensor (avr.) · Fleet (avr.) · Ada-MK (mai)
   │    AutoMegaKernel (juin)
   │    Dynamisme, chiplets, production, synthèse automatique
   ▼

2.4 Les persistent threads, la source

Le papier fondateur est celui de Gupta, Stuart et Owens (UC Davis, 2012), A Study of Persistent Threads Style GPU Programming for GPGPU Workloads.

Sa contribution : caractériser formellement un style de programmation où l'on « circonscrit le noyau logique dans une boucle, de sorte que cette boucle continue de tourner tant qu'il reste des éléments à traiter ».

__global__ void persistant(FileTravail* file) {
    while (true) {
        int tache = file->prochaine();      // atomicAdd sur un compteur
        if (tache < 0) break;               // plus de travail
        traiter(tache);
    }
}

Les auteurs identifient quatre cas d'usage, et il est frappant de constater que les megakernels de 2025-2026 les exploitent tous les quatre simultanément :

Cas d'usage (2012) Dans un megakernel (2026)
Synchronisation CPU-GPU le CPU pousse des instructions, le GPU les consomme
Équilibrage de charge file de travail dynamique entre SM
Localité producteur-consommateur activations gardées en mémoire partagée
Synchronisation globale compteurs de dépendances

Leur conclusion mérite d'être citée : 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 ».

C'est-à-dire : la technique est puissante et pas universellement gagnante. C'est exactement le message du chapitre 9, et il n'a pas changé en quatorze ans.

2.5 L'anatomie générale

Tous les megakernels publiés partagent la même structure, avec des variations de vocabulaire.

┌──────────────────────────────────────────────────────────────────┐
│                        MEGAKERNEL                                │
│                                                                  │
│  CPU : construit la liste d'instructions / le graphe de tâches   │
│           │                                                      │
│           ▼                                                      │
│  ┌────────────────────────────────────────────────────────────┐  │
│  │ SM 0            SM 1            …            SM 147        │  │
│  │ ┌────────┐      ┌────────┐                  ┌────────┐     │  │
│  │ │loader  │      │loader  │                  │loader  │     │  │
│  │ │compute │      │compute │                  │compute │     │  │
│  │ │storer  │      │storer  │                  │storer  │     │  │
│  │ └────────┘      └────────┘                  └────────┘     │  │
│  │    │                │                           │          │  │
│  │    └────────────────┴───────────────────────────┘          │  │
│  │              mémoire partagée paginée                       │  │
│  └────────────────────────────────────────────────────────────┘  │
│           │                                                      │
│           ▼                                                      │
│  Mémoire globale : compteurs de dépendances, files de tâches,    │
│                     activations intermédiaires                   │
└──────────────────────────────────────────────────────────────────┘

Le vocabulaire selon les travaux :

Concept Hazy Research Mirage MPK Event Tensor
Unité de travail instruction tâche (task) tâche tuilée
Synchronisation compteurs événements (events) event tensor
Représentation liste par SM ttGraph (SM-level graph) tenseur d'événements
Répartition interpréteur + file workers + schedulers statique + dynamique
Mémoire partagée pages de 16 Ko pages de 32 Ko paginée

2.6 Les cinq mécanismes indispensables

Quel que soit le travail, cinq briques sont nécessaires.

1. La persistance

Une grille dimensionnée pour la résidence complète :

int blocs_par_sm;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocs_par_sm, megakernel,
                                              THREADS, SMEM);
int nb_sm;
cudaDeviceGetAttribute(&nb_sm, cudaDevAttrMultiProcessorCount, 0);
megakernel<<<nb_sm * blocs_par_sm, THREADS, SMEM>>>(...);

Dépasser cette taille provoque un interblocage.

2. Le dispatch d'instructions

while (true) {
    Instruction ins = prochaine_instruction();
    if (ins.opcode == FIN) break;
    switch (ins.opcode) {
        case RMSNORM_QKV_ROPE: executer<RmsNormQkvRope>(ins); break;
        case ATTENTION:        executer<Attention>(ins);      break;
        case O_PROJ:           executer<OProj>(ins);          break;
        // ...
    }
}

Le switch divergent, mais pas coûteux

Ce switch porte sur un opcode uniforme dans tout le bloc : tous les threads d'un SM exécutent la même instruction. Il n'y a donc aucune divergence intra-warp.

L'hétérogénéité est entre SM, ce qui est gratuit. C'est le point souligné en Domaines 6, §6.4 : la divergence bien organisée ne coûte rien.

3. Les compteurs de dépendance

// Producteur, à la fin de son travail
__threadfence();                                  // rendre les données visibles
if (thread_representant) {
    atomicAdd(&compteurs[ins.compteur_sortie], 1);
}

// Consommateur, avant de commencer
if (thread_representant) {
    while (atomicAdd(&compteurs[ins.compteur_entree], 0) < ins.valeur_cible) {
        __nanosleep(20);                          // réduire la contention
    }
}
__syncthreads();

4. L'allocateur de mémoire partagée

La mémoire partagée du SM devient un tas géré par pages, puisqu'un seul bloc persistant l'occupe pendant toute l'exécution et que les instructions successives ont des besoins différents.

5. L'ordonnancement

Statique (chaque SM a sa liste), dynamique (file de travail globale), ou hybride.

2.7 Le lien avec la spécialisation des warps

Un megakernel est, structurellement, la spécialisation des warps portée à l'échelle du GPU entier.

Concept Dans une GEMM Hopper Dans un megakernel
Unité de travail une tuile une instruction
Producteur 1 warp group émettant TMA threads loader
Consommateur 2 warp groups wgmma threads compute
Épilogue dans le consommateur threads storer
Tampon mémoire partagée circulaire pages de mémoire partagée
Synchronisation mbarrier (locale au bloc) compteurs globaux
Ordonnancement statique file de travail + SM ordonnanceurs

C'est pourquoi la partie 4 était un prérequis : un megakernel est un noyau warp-specialized dont l'horizon est le modèle entier plutôt qu'une tuile.

Résumé du chapitre

À retenir

  • Un megakernel : un seul lancement, persistant, avec des dépendances fines et une hétérogénéité interne.
  • Ce n'est ni de la fusion simple (qui reste dans un bloc), ni un CUDA Graph (qui garde les frontières), ni un noyau persistant ordinaire (qui n'a qu'une opération).
  • L'idée vient des persistent threads de Gupta, Stuart et Owens (2012), dont les quatre cas d'usage sont tous exploités simultanément par les megakernels modernes.
  • Cinq mécanismes indispensables : persistance, dispatch, compteurs, allocateur de mémoire partagée, ordonnancement.
  • Le switch d'opcode ne coûte rien : il est uniforme dans le bloc, l'hétérogénéité est entre SM.
  • Un megakernel est la spécialisation des warps à l'échelle du GPU.

Vérifiez que vous avez compris

Pourquoi la persistance est-elle obligatoire, et pas seulement souhaitable ?

Parce que la synchronisation inter-blocs par compteurs interbloque si tous les blocs ne sont pas résidents.

Scénario : le bloc A attend le compteur que le bloc B doit incrémenter. Si B n'a pas encore été lancé — parce qu'aucun SM ne s'est libéré, et qu'aucun ne se libérera puisque A attend — le programme se fige.

C'est pourquoi tout megakernel dimensionne sa grille par cudaOccupancyMaxActiveBlocksPerMultiprocessor × nb_SM, et pourquoi cudaLaunchCooperativeKernel impose la même contrainte.

En quoi le mécanisme de MPK diffère-t-il de simples compteurs ?

MPK introduit une notion d'événement avec des propriétés structurelles fortes.

Dans son graphe (le ttGraph), « tâches et événements alternent : chaque tâche n'a que des arêtes sortantes vers des événements déclencheurs et des arêtes entrantes depuis des événements de dépendance ». Après normalisation, chaque tâche a au plus un événement de dépendance et un événement déclencheur.

Cette régularité permet au compilateur d'appliquer des optimisations — fusion d'événements par ensembles de successeurs et de prédécesseurs, linéarisation par parcours en largeur pour que les tâches déclenchées par le même événement occupent des indices contigus.

Autrement dit : ce sont des compteurs, avec une structure algébrique qui les rend optimisables par un compilateur. C'est la différence entre écrire un megakernel à la main et en générer un.

Le papier de 2012 dit que les persistent threads « peuvent aussi entraîner une perte de performance dans de nombreux cas ». Lesquels ?

Quatre catégories, toutes toujours valables :

  1. Quand le travail est régulier et abondant. L'ordonnanceur matériel équilibre déjà parfaitement ; la file logicielle n'ajoute que du coût atomique.
  2. Quand l'occupancy en souffre. Un noyau persistant consomme des registres et de la mémoire partagée pour son état d'ordonnancement, ce qui réduit le nombre de warps résidents.
  3. Quand la contention atomique domine. Avec de nombreux blocs sondant la même file, les atomiques se sérialisent.
  4. Quand le noyau devient trop gros. Pression sur les registres, spilling, temps de compilation, et perte de la spécialisation par opération.

Ces quatre points sont exactement ceux du chapitre 9.

Chapitre suivant : 3 · L'interpréteur sur GPU

Sources de ce chapitre