Aller au contenu

1 · Le problème des frontières de noyau

Ce que coûte, exactement, le fait de terminer un noyau et d'en lancer un autre. Trois coûts distincts, souvent confondus.

1.1 Le coût 1 : le lancement

C'est le plus connu et le moins important.

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

Ce coût couvre : préparation des arguments côté CPU, appel au pilote, écriture dans le tampon de commandes, signalisation au GPU, et configuration des ressources (mémoire partagée, registres) sur les SM.

CUDA Graphs le divise par 4 à 8, en pré-enregistrant toute la séquence. C'est pourquoi vLLM, SGLang, TensorRT-LLM et torch.compile(mode="reduce-overhead") l'utilisent tous.

Et c'est pourquoi le lancement n'est pas le vrai problème : à 1,3 µs par noyau et 224 noyaux pour un modèle de 32 couches, on est à 291 µs. Significatif, mais ce n'est qu'une partie.

1.2 Le coût 2 : la barrière globale implicite

Voici le coût que CUDA Graphs ne supprime pas.

La règle du modèle CUDA : aucun bloc du noyau \(n+1\) ne peut démarrer avant que tous les blocs du noyau \(n\) soient terminés.

Considérons un noyau à 512 blocs sur un GPU à 148 SM :

vague 1 : blocs   0-147   ████████████
vague 2 : blocs 148-295   ████████████
vague 3 : blocs 296-443   ████████████
vague 4 : blocs 444-511   ████░░░░░░░░   ← 68 blocs seulement
                                          80 SM inactifs
                          ↑
                    puis barrière globale : personne ne démarre
                    le noyau suivant avant la fin de la vague 4

Hazy Research le formule ainsi : « avec 512 blocs mais seulement 148 multiprocesseurs, 80 SM restent inactifs à attendre les retardataires ».

Deux effets se cumulent :

  1. la queue de vague : la dernière vague est partielle, une fraction du GPU ne fait rien ;
  2. la barrière : même les SM qui ont fini ne peuvent pas commencer le travail suivant.

Ce coût est structurel. Il ne dépend pas du pilote ni du CPU : il est dans le modèle d'exécution.

1.3 Le coût 3 : la bulle mémoire

Le plus insidieux, et probablement le plus coûteux.

Au moment où un noyau démarre, aucune requête mémoire n'est en vol. Il faut donc attendre plusieurs centaines de cycles avant que les premières données arrivent. Pendant ce temps, les unités de calcul ne font rien, et surtout : la bande passante mémoire n'est pas utilisée.

Un noyau isolé :
  temps →
  [émission des premières requêtes]
  [......attente ~500 cycles......]
  [calcul ∥ chargements suivants]
  [drainage]
  ← puis fin de noyau, tout s'arrête

Bande passante utilisée :
  ░░░░░░░░░░░░████████████████░░░░░░
       bulle      régime établi   drainage

Sur une séquence de noyaux courts, ces bulles représentent une fraction importante du temps. Et surtout : elles empêchent la mémoire d'être continûment sollicitée, ce qui est exactement ce qu'il faut faire quand on est limité par la bande passante.

L'analyse chiffrée de Hazy Research sur Llama-1B :

  • la bande passante d'un H100 permettrait ~1 350 passes avant par seconde ;
  • les systèmes réels en réalisent ~770 ;
  • soit 57 % de la performance théorique, avec « environ cinq microsecondes de blocage par noyau, sur 7 lancements par couche et 16 couches ».

1.4 Le décompte complet

Le tableau qui résume, pour une passe avant de Llama-1B sur B200 (600 µs au total) :

Poste Temps Nature
Stockage des activations, attente de cohérence, rechargement 250 µs frontières
RMSNorm + produits matrice-vecteur (95 % pour les matvec) 200 µs calcul utile
Attente du chargement des poids 30 µs mémoire
Synchronisation entre warps (~60 ns × N) 40 µs frontières
Configuration et divers 80 µs frontières

Total « frontières » : 370 µs sur 600, soit 62 %.

Le chiffre à retenir de ce chapitre

Sur une passe avant de modèle de langage à petit lot, plus de la moitié du temps n'est pas du calcul : c'est du stockage, de la synchronisation, de l'attente et de la configuration.

Aucune optimisation de noyau individuel ne peut récupérer cela. Il faut changer la structure d'exécution.

1.5 Pourquoi CUDA Graphs ne suffit pas

Récapitulons ce que CUDA Graphs supprime et ne supprime pas.

Coût CUDA Graphs
Lancement (CPU) divisé par 4-8
Barrière globale implicite inchangée
Bulle mémoire inchangée
Allers-retours d'activations en HBM inchangés

CUDA Graphs élimine le surcoût CPU. Le reste est dans le modèle d'exécution GPU.

C'est un point souvent mal compris : le chiffre de « 1,3 µs par lancement » cité par Hazy Research est déjà avec CUDA Graphs. Ils partent d'une ligne de base optimisée.

1.6 Pourquoi PDL ne suffit pas non plus

Hopper introduit le Programmatic Dependent Launch : le noyau \(n+1\) peut exécuter son prologue (chargement de ses poids) pendant que le noyau \(n\) termine.

__global__ void suivant() {
    charger_poids();                    // avant la dépendance
    cudaGridDependencySynchronize();    // ← attendre le noyau précédent
    calculer();
}

C'est un vrai progrès sur la bulle mémoire. Mais la granularité est trop grossière.

L'exemple donné par Hazy Research : dans un MLP, la projection descendante (down_proj) a besoin de l'état caché produit par up/gate. Avec PDL, il faut attendre que tout l'état caché soit produit.

Or down_proj pourrait commencer dès que le premier quart est prêt.

Leur solution : découper en quatre morceaux avec un compteur par morceau. Chaque morceau débloque son consommateur indépendamment.

PDL :
  [up/gate : 4 morceaux .......................]
                                               [down : 4 morceaux ......]

Compteurs fins :
  [up/gate m0][m1][m2][m3]
              [down m0][m1][m2][m3]
              ↑ démarre dès que m0 est prêt

Ce découpage n'est pas exprimable avec PDL. Il l'est naturellement dans un megakernel.

1.7 La hiérarchie des solutions

┌────────────────────────────────────────────────────────────────┐
│ Coût d'une frontière                    Ce qui est supprimé    │
├────────────────────────────────────────────────────────────────┤
│ Lancement naïf         ~5-10 µs         —                      │
│  ↓                                                             │
│ CUDA Graphs            ~1,3 µs          surcoût CPU            │
│  ↓                                                             │
│ PDL                    ~1,3 µs          une partie de la bulle │
│  ↓                                                             │
│ Cooperative groups     µs par barrière  la frontière, pas le   │
│ (grid.sync)                             coût de synchronisation│
│  ↓                                                             │
│ MEGAKERNEL             ~0,1 µs          tout, sauf la          │
│ (compteurs fins)                        dépendance réelle      │
└────────────────────────────────────────────────────────────────┘

La dernière ligne mérite d'être explicitée. Un megakernel ne supprime pas les dépendances — elles sont dans le modèle. Il supprime tout ce qui n'est pas la dépendance elle-même : le surcoût de lancement, la barrière sur des blocs sans rapport, la bulle mémoire, et l'aller-retour en HBM des activations.

1.8 Le calcul de rentabilité

Quand ce raisonnement s'applique-t-il ?

Notons \(t_{\text{calcul}}\) le temps de calcul utile d'une passe et \(t_{\text{frontières}}\) le surcoût.

\[ \text{gain potentiel} = \frac{t_{\text{frontières}}}{t_{\text{calcul}} + t_{\text{frontières}}} \]
Situation \(t_{\text{calcul}}\) \(t_{\text{frontières}}\) Gain potentiel
Entraînement, gros lot 500 ms 0,3 ms 0,06 %
Préremplissage, invite longue 50 ms 0,3 ms 0,6 %
Décodage, 70 B, lot 64 25 ms 0,4 ms 1,6 %
Décodage, 8 B, lot 8 6 ms 0,3 ms 5 %
Décodage, 1 B, lot 1 0,6 ms 0,4 ms 40 %

La conclusion du chapitre

Les megakernels sont rentables exactement là où le travail par passe est faible : petits modèles, petits lots, décodage.

C'est un domaine étroit, et c'est aussi le domaine qui compte le plus en 2026 : les agents, la génération de code interactive, les modèles embarqués, et tout ce qui est latence-critique.

Ce raisonnement est repris et détaillé en chapitre 9.

Résumé du chapitre

À retenir

  • Une frontière de noyau coûte trois choses : le lancement (~1,3 µs avec CUDA Graphs), la barrière globale implicite, et la bulle mémoire.
  • CUDA Graphs ne supprime que la première.
  • PDL réduit la bulle mais sa granularité est trop grossière : il faut attendre la fin complète du noyau précédent.
  • Sur Llama-1B/B200, 62 % du temps d'une passe avant est du stockage, de la synchronisation, de l'attente et de la configuration — pas du calcul.
  • Les systèmes réels atteignent ~57 % de ce que la bande passante permettrait.
  • Le gain potentiel d'un megakernel est proportionnel à la part du surcoût : ~40 % à lot 1 sur un petit modèle, < 1 % en entraînement.

Vérifiez que vous avez compris

Pourquoi la barrière globale implicite coûte-t-elle même quand tous les blocs durent le même temps ?

Parce que le problème n'est pas la variance mais la quantification de vagues.

Avec 512 blocs et 148 SM, il y a 3,46 vagues. La dernière contient 68 blocs au lieu de 148 : 80 SM sont inactifs pendant toute sa durée, même si chaque bloc dure exactement le même temps.

Perte : \(80/148 \times 1/3{,}46 \approx 15\ \%\) du temps de ce noyau.

À cela s'ajoute, s'il y a de la variance, l'attente du bloc le plus lent. Les deux effets se cumulent.

Un modèle de 70 milliards de paramètres à lot 1 sur 8 GPU. Un megakernel est-il rentable ?

Calculons.

  • Poids par GPU : \(140/8 = 17{,}5\) Go ;
  • \(t_{\min} = 17{,}5 / 3{,}35 = 5{,}2\) ms par jeton ;
  • couches : 80, soit ~7 × 80 = 560 noyaux ;
  • surcoût de lancement avec CUDA Graphs : \(560 \times 1{,}3 = 728\ \mu s\) ;
  • plus les barrières, les bulles et les allers-retours d'activations, qu'on peut estimer du même ordre voire davantage.

Surcoût total plausible : 1,5 à 2,5 ms sur ~7 ms de temps réel, soit 20 à 35 %.

C'est significatif — et c'est cohérent avec les 1,1-1,4× que MPK annonce sur 8 H100 en tensor-parallèle. À cela s'ajoute le bénéfice du recouvrement calcul/communication, qui n'est pas dans ce décompte.

Pourquoi ne pas simplement fusionner tous les noyaux d'une couche en un seul gros noyau, sans construire un megakernel complet ?

C'est possible, et cela ne suffit pas, pour deux raisons.

1. La synchronisation inter-blocs reste nécessaire à l'intérieur. Une couche contient des dépendances qu'un bloc seul ne peut pas satisfaire : la projection de sortie a besoin du résultat de l'attention de toutes les têtes, calculées par des blocs différents. Fusionner exige donc déjà le mécanisme de compteurs — c'est-à-dire l'essentiel du travail d'un megakernel.

2. Le gain est proportionnel au nombre de frontières supprimées. Fusionner par couche supprime ~6 frontières sur 7 par couche, mais en laisse une entre couches : 32 frontières restantes au lieu de 224. C'est déjà beaucoup mieux.

En pratique, une fois qu'on a payé le prix du mécanisme, il n'y a plus de raison de s'arrêter à la couche. D'où le nom : megakernel.

Chapitre suivant : 2 · Définition et généalogie

Sources de ce chapitre