Aller au contenu

6 · Flux, événements et graphes

Le parallélisme au-dessus du noyau. Ce chapitre mène directement aux megakernels : il explique ce qu'on peut faire pour réduire le coût des frontières de noyau sans les supprimer, et où cette approche atteint sa limite.

6.1 Les flux

Un flux (stream) est une file d'opérations exécutées dans l'ordre. Deux flux différents peuvent s'exécuter dans n'importe quel ordre relatif, et donc en parallèle.

cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);

noyau_a<<<g, b, 0, s1>>>(...);     // ces deux noyaux peuvent
noyau_b<<<g, b, 0, s2>>>(...);     // s'exécuter simultanément

cudaStreamSynchronize(s1);
cudaStreamSynchronize(s2);

Sans flux explicite, tout va dans le flux nul (default stream), qui est sérialisé.

Le flux nul a une sémantique spéciale

Par défaut, le flux nul se synchronise implicitement avec tous les autres flux : une opération dans le flux nul attend que tous les autres flux soient vides, et les bloque ensuite. C'est une source de sérialisation invisible.

Depuis CUDA 7, on peut compiler avec --default-stream per-thread pour que chaque thread CPU ait son propre flux non bloquant. C'est presque toujours ce qu'on veut.

Créez vos flux avec cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking) pour éviter la synchronisation implicite avec le flux nul.

À quoi ça sert vraiment

Trois usages, par ordre d'importance.

1. Recouvrir les transferts et le calcul.

const int n_morceaux = 4;
for (int i = 0; i < n_morceaux; ++i) {
    int off = i * taille_morceau;
    cudaMemcpyAsync(d + off, h + off, octets_morceau,
                    cudaMemcpyHostToDevice, flux[i]);
    noyau<<<g, b, 0, flux[i]>>>(d + off, taille_morceau);
    cudaMemcpyAsync(h + off, d + off, octets_morceau,
                    cudaMemcpyDeviceToHost, flux[i]);
}

sans flux :  [H→D][calcul][D→H][H→D][calcul][D→H]…
avec flux :  [H→D][H→D][H→D][H→D]
                  [calcul][calcul][calcul][calcul]
                          [D→H][D→H][D→H][D→H]

Gain théorique : jusqu'à 3× si les trois phases durent autant. Condition obligatoire : la mémoire hôte doit être épinglée (cudaMallocHost), sinon cudaMemcpyAsync est en réalité synchrone.

2. Occuper le GPU avec plusieurs petits noyaux. Si un noyau n'utilise que 30 SM sur 132, deux autres noyaux peuvent tourner en même temps. Le matériel les placera sur des SM libres.

3. Isoler des travaux concurrents — par exemple plusieurs requêtes d'inférence.

6.2 Les événements

Un événement est un marqueur inséré dans un flux.

Usage 1 : mesurer le temps correctement

cudaEvent_t debut, fin;
cudaEventCreate(&debut);
cudaEventCreate(&fin);

// Échauffement — indispensable
for (int i = 0; i < 10; ++i) noyau<<<g, b>>>(...);
cudaDeviceSynchronize();

cudaEventRecord(debut);
for (int i = 0; i < 100; ++i) noyau<<<g, b>>>(...);
cudaEventRecord(fin);
cudaEventSynchronize(fin);

float ms;
cudaEventElapsedTime(&ms, debut, fin);
printf("%.3f µs par lancement\n", ms * 1000.0f / 100.0f);

Les événements CUDA sont horodatés par le GPU, ce qui élimine le bruit du CPU et de la mise en file.

L'échauffement n'est pas optionnel

La première exécution d'un noyau inclut : la compilation JIT du PTX si nécessaire (des dizaines de millisecondes), l'allocation du contexte, la pagination de la mémoire, et une fréquence GPU encore basse. Mesurer sans échauffement donne des chiffres faux d'un ordre de grandeur.

Pour des mesures reproductibles, verrouillez aussi les fréquences :

sudo nvidia-smi -lgc 1410,1410     # verrouille l'horloge graphique
sudo nvidia-smi --lock-memory-clocks=1593

Usage 2 : dépendances entre flux

noyau_a<<<g, b, 0, s1>>>(...);
cudaEventRecord(evt, s1);
cudaStreamWaitEvent(s2, evt, 0);    // s2 attend que noyau_a soit fini
noyau_b<<<g, b, 0, s2>>>(...);      // ne démarre qu'après noyau_a

C'est ainsi qu'on construit un graphe de dépendances arbitraire entre noyaux.

6.3 Le coût d'un lancement de noyau

Voici le chiffre autour duquel tourne toute la partie 8.

Situation Coût par lancement
Lancement standard ~5 à 10 µs (CPU)
Lancement dans un graphe CUDA ~1,3 µs
Mesuré sur MI300X (HIP) ~4,5 µs

Les 5 à 10 µs du lancement standard se décomposent en : coût CPU de préparation des arguments, appel au pilote, écriture dans le tampon de commandes, et signalisation au GPU.

Pourquoi c'est un problème. Une passe avant de Llama-1B, c'est ~100 noyaux. À 5 µs chacun, cela fait 500 µs de pur surcoût — alors que la passe avant elle-même dure ~1 ms. Le surcoût représente un tiers du temps total.

À cela s'ajoutent deux coûts que le lancement seul ne mesure pas :

  1. la barrière globale implicite. Aucun bloc du noyau \(n+1\) ne démarre avant que tous les blocs du noyau \(n\) soient terminés. Hazy Research illustre : « avec 512 blocs et seulement 148 SM, 80 SM restent inactifs à attendre les retardataires » ;
  2. la bulle mémoire. Au démarrage d'un noyau, aucune requête mémoire n'est en vol. Il faut de nouveau attendre des centaines de cycles.

L'analyse chiffrée de Hazy Research : la bande passante d'un H100 permettrait ~1 350 passes avant par seconde sur Llama-1B ; les systèmes réels en font ~770, avec « cinq microsecondes de blocage par noyau, sur 7 lancements par couche et 16 couches ».

6.4 CUDA Graphs

L'idée : au lieu de soumettre 100 noyaux un par un, on enregistre la séquence une fois et on la rejoue.

cudaGraph_t graphe;
cudaGraphExec_t graphe_exec;

// 1. Capture
cudaStreamBeginCapture(flux, cudaStreamCaptureModeGlobal);
for (int i = 0; i < 100; ++i) {
    noyau_i<<<g, b, 0, flux>>>(...);
}
cudaStreamEndCapture(flux, &graphe);

// 2. Instanciation (une seule fois)
cudaGraphInstantiate(&graphe_exec, graphe, nullptr, nullptr, 0);

// 3. Exécution (à chaque itération)
cudaGraphLaunch(graphe_exec, flux);

Ce que ça achète :

  • le coût CPU par noyau s'effondre : un seul appel au pilote pour tout le graphe. Le lancement passe de ~5-10 µs à ~1,3 µs par nœud ;
  • le pilote peut pré-calculer les dépendances et pré-allouer les ressources ;
  • les dépendances explicites du graphe permettent au matériel de démarrer certains nœuds en avance.

Ce que ça n'achète pas :

  • la barrière globale entre noyaux reste. Les nœuds dépendants s'attendent toujours ;
  • la bulle mémoire reste ;
  • les formes et les pointeurs sont figés dans le graphe. Changer une taille de lot exige de recapturer (ou d'utiliser cudaGraphExecUpdate).

Ce qu'il faut retenir de CUDA Graphs

C'est la première étape vers l'élimination du surcoût de lancement, et elle est presque gratuite : torch.cuda.graphs, torch.compile(mode="reduce-overhead"), vLLM et SGLang l'utilisent tous par défaut.

Et elle ne suffit pas. Le chiffre de 1,3 µs par lancement cité par Hazy Research est déjà avec CUDA Graphs. C'est précisément parce que CUDA Graphs ne supprime ni la barrière ni la bulle que les megakernels existent.

6.5 Programmatic Dependent Launch

Introduit avec Hopper, PDL permet à un noyau de commencer son prologue (chargement de ses poids) avant que le noyau précédent soit totalement terminé.

__global__ void noyau_suivant(...) {
    // Prologue : charger ce qui ne dépend pas du noyau précédent
    charger_poids();

    cudaGridDependencySynchronize();   // ← attendre le noyau précédent ici

    // Corps : utiliser les résultats du noyau précédent
    calculer();
}

Et dans le noyau précédent :

__global__ void noyau_precedent(...) {
    calculer();
    cudaTriggerProgrammaticLaunchCompletion();   // libère le suivant plus tôt
}

C'est un vrai progrès, et il reste trop grossier. Hazy Research le dit explicitement : PDL force à attendre la fin complète du noyau précédent avant d'utiliser le moindre de ses résultats. Or dans un transformeur, la projection descendante du MLP pourrait commencer dès que le premier quart de l'état caché est prêt.

Leur solution : découper en quatre morceaux avec un compteur par morceau. Ce qui n'est pas exprimable avec PDL — et qui l'est naturellement dans un megakernel.

6.6 L'échelle du problème, résumée

┌────────────────────────────────────────────────────────────┐
│ Coût d'une frontière de noyau                              │
├────────────────────────────────────────────────────────────┤
│ Lancement naïf                     ~5-10 µs (CPU)          │
│  ↓ CUDA Graphs                                             │
│ Lancement dans un graphe            ~1,3 µs                │
│  ↓ PDL                                                     │
│ Prologue recouvert                  ~1,3 µs, bulle réduite │
│  ↓ Megakernel                                              │
│ Dépendance par compteur             ~0,1 µs, aucune bulle  │
└────────────────────────────────────────────────────────────┘

Chaque ligne divise le coût. La dernière exige de réécrire complètement la façon dont on structure un modèle, et c'est le sujet de la partie 8.

6.7 Les Green Contexts

Une nouveauté de CUDA 13.1 qui mérite d'être connue : les Green Contexts permettent de partitionner un GPU au niveau des SM, avec une API runtime.

// Créer un contexte disposant de 64 SM sur 132
cudaGreenCtx_t ctx;
// (API simplifiée ; voir la documentation pour la forme exacte)

Usage : garantir qu'une tâche latence-critique dispose de ses SM sans être perturbée par une tâche de fond. C'est une alternative plus fine que MPS (Multi-Process Service) ou MIG (Multi-Instance GPU), qui partitionnent au niveau du processus.

Pour les megakernels, c'est intéressant conceptuellement : un megakernel s'approprie déjà tout le GPU. Les Green Contexts offrent un moyen de lui en donner seulement une partie.

Résumé du chapitre

À retenir

  • Un flux est une file ordonnée ; deux flux s'exécutent en parallèle. Créez vos flux avec cudaStreamNonBlocking pour éviter la sémantique spéciale du flux nul.
  • Recouvrir transferts et calcul exige de la mémoire hôte épinglée.
  • Mesurez avec des événements CUDA, après échauffement, sur plusieurs répétitions, à fréquences verrouillées.
  • Un lancement de noyau coûte 5 à 10 µs, ramené à ~1,3 µs par CUDA Graphs. C'est le plancher de l'approche « un noyau par opérateur ».
  • CUDA Graphs ne supprime ni la barrière globale ni la bulle mémoire. PDL les réduit mais reste trop grossier.
  • D'où les megakernels.

Vérifiez que vous avez compris

Vous utilisez cudaMemcpyAsync avec de la mémoire allouée par malloc. Le recouvrement ne fonctionne pas. Pourquoi ?

Parce que cudaMemcpyAsync n'est asynchrone que depuis de la mémoire épinglée. Avec de la mémoire paginable, le pilote doit d'abord copier vers un tampon interne épinglé — opération synchrone qui bloque le CPU.

Le remède : cudaMallocHost (ou cudaHostRegister sur une allocation existante). Gain secondaire : le débit PCIe passe de ~6 Go/s à ~55 Go/s.

Un modèle a 100 noyaux et vous activez CUDA Graphs. Le temps passe de 1,5 ms à 1,1 ms. Cohérent ?

Oui. 100 noyaux × (5 µs − 1,3 µs) ≈ 370 µs économisés, ce qui correspond aux 400 µs observés. Le résidu (100 × 1,3 = 130 µs de lancements + les barrières + les bulles mémoire) reste dans les 1,1 ms.

C'est exactement l'ordre de grandeur qui rend un megakernel intéressant : il reste ~10 à 30 % à gagner, et sur du décodage à lot 1 c'est énorme.

Pourquoi torch.compile(mode=\"reduce-overhead\") échoue-t-il quand la taille du lot change à chaque appel ?

Ce mode active CUDA Graphs, et un graphe capture des pointeurs et des formes figés. Un changement de forme invalide le graphe et déclenche une recapture, coûteuse.

C'est le même problème que rencontre Mirage MPK, qui doit générer un ttGraph par taille de lot représentative, et que le papier Event Tensor attaque de front en introduisant du dynamisme dans la représentation.

Chapitre suivant : 7 · Atomiques et synchronisation

Sources de ce chapitre