1 · Le modèle roofline¶

Le meilleur rapport entre effort intellectuel et information obtenue de toute la programmation GPU. Cinq minutes de papier valent parfois une journée de profileur.

1.1 L'idée¶

Un noyau consomme deux ressources : du calcul et de la bande passante mémoire. Il ne peut pas aller plus vite que la plus contrainte des deux.

Notons :

Symbole	Signification	Unité
\(W\)	travail total du noyau	opérations (FLOP)
\(Q\)	données transférées depuis/vers la mémoire globale	octets
\(P\)	débit de calcul crête de la machine	opérations/s
\(B\)	bande passante mémoire crête de la machine	octets/s
\(I\)	intensité arithmétique du noyau, \(I = W/Q\)	opérations/octet

Le temps d'exécution est borné par :

\[ t \ge \max\left( \frac{W}{P},\ \frac{Q}{B} \right) \]

et la performance atteignable est :

\[ \text{perf} = \frac{W}{t} \le \min\left( P,\ I \times B \right) \]

C'est l'équation du roofline. Tracée en échelle logarithmique, elle donne une courbe en forme de toit :

performance
(FLOP/s)
    │                    ┌──────────────────────  P (plafond de calcul)
    │                   ╱
    │                  ╱
    │                 ╱   pente = B (plafond mémoire)
    │                ╱
    │               ╱
    │              ╱
    └─────────────┴───────────────────────────→  intensité arithmétique I
                  I_crit = P/B                    (FLOP/octet)

Le point d'inflexion, appelé intensité critique ou ridge point :

\[ I_{\text{crit}} = \frac{P}{B} \]

\(I < I_{\text{crit}}\) → limité par la mémoire (memory-bound) ;
\(I > I_{\text{crit}}\) → limité par le calcul (compute-bound).

1.2 Les intensités critiques réelles¶

Le chiffre dépend de la précision utilisée, ce qui est essentiel.

Machine	Précision	\(P\)	\(B\)	\(I_{\text{crit}}\)
A100	FP32	19,5 TFLOPS	2,0 To/s	~10
A100	BF16 tensor	312 TFLOPS	2,0 To/s	~156
H100	FP32	67 TFLOPS	3,35 To/s	~20
H100	BF16 tensor	990 TFLOPS	3,35 To/s	~296
H100	FP8 tensor	1 979 TFLOPS	3,35 To/s	~591
B200	BF16 tensor	~2 250 TFLOPS	8,0 To/s	~281
B200	FP4 tensor	~9 000 TFLOPS	8,0 To/s	~1 125
MI300X	BF16	1 307 TFLOPS	5,3 To/s	~247

La conséquence stratégique du tableau

Chaque génération augmente \(P\) plus vite que \(B\). L'intensité critique monte, ce qui veut dire que de plus en plus de noyaux basculent du côté « limité par la mémoire ».

Sur un H100 en BF16, il faut 296 opérations par octet lu pour saturer les tensor cores. C'est énorme. Une GEMM y arrive (une GEMM \(N^3\) sur des matrices \(N^2\) a \(I \propto N\)), presque rien d'autre.

C'est le fait structurant de la programmation GPU en 2026, et la raison d'être de la fusion de noyaux, de la quantification, et des megakernels.

1.3 Trois calculs à faire de tête¶

Exemple 1 — addition élémentaire¶

c[i] = a[i] + b[i];    // float

\(W = 1\) opération par élément ;
\(Q = 12\) octets (2 lectures + 1 écriture de 4 octets) ;
\(I = 1/12 \approx 0{,}083\).

Sur H100, \(I_{\text{crit}} = 20\). On est 240 fois en dessous. Ce noyau est irrémédiablement limité par la mémoire, et sa performance maximale est :

\[ \text{perf} = 0{,}083 \times 3{,}35 \times 10^{12} = 0{,}28 \text{ TFLOPS} \]

soit 0,4 % du pic FP32. Ce n'est pas un échec : c'est la nature du problème.

Le seul indicateur pertinent devient la bande passante atteinte, pas les FLOPS :

\[ t_{\min} = \frac{12n}{3{,}35 \times 10^{12}} \]

Pour \(n = 10^8\) : \(t_{\min} = 358\ \mu s\).

La bonne métrique dépend du régime

Noyau limité par la mémoire → mesurez en Go/s et comparez à \(B\).
Noyau limité par le calcul → mesurez en TFLOPS et comparez à \(P\).

Annoncer « mon noyau d'addition fait 0,3 TFLOPS » n'a aucun sens. Annoncer « il atteint 3,1 To/s sur 3,35 possibles, soit 93 % » est une information.

Exemple 2 — GEMM¶

Pour \(\mathbf{C}_{M \times N} = \mathbf{A}_{M \times K} \mathbf{B}_{K \times N}\) :

\(W = 2MNK\) ;
\(Q = 2(MK + KN + MN)\) octets en BF16, si chaque élément n'est lu qu'une fois ;
pour \(M = N = K = 4096\) : \(W = 1{,}37 \times 10^{11}\), \(Q = 1{,}01 \times 10^{8}\), \(I \approx 1\,365\).

\(I = 1365 \gg I_{\text{crit}} = 296\) : la GEMM est limitée par le calcul, et c'est pour cela que les tensor cores existent.

\[ t_{\min} = \frac{1{,}37 \times 10^{11}}{990 \times 10^{12}} = 139\ \mu s \]

Attention : ce \(I\) suppose une réutilisation parfaite. Un noyau naïf a \(I = 1/4\) et se retrouve du mauvais côté. L'intensité arithmétique n'est pas une propriété du problème, c'est une propriété de l'implémentation. C'est exactement ce que fait le pavage en tuiles : il déplace le noyau vers la droite du graphique.

Exemple 3 — décodage LLM à lot 1¶

C'est le calcul qui explique toute la partie 8.

Pour un modèle de \(\Theta\) paramètres en BF16, générer un jeton exige de lire tous les poids :

\(Q \approx 2\Theta\) octets ;
\(W \approx 2\Theta\) opérations (chaque poids sert à une multiplication-accumulation) ;
\(I \approx 1\).

Une opération par octet. Contre un seuil de 296. On est 296 fois du mauvais côté, et aucune optimisation de calcul n'y changera rien.

Le temps minimal par jeton pour Llama-70B en BF16 sur un H100 :

\[ t_{\min} = \frac{2 \times 70 \times 10^9}{3{,}35 \times 10^{12}} = 41{,}8\ \text{ms} \]

soit 24 jetons par seconde au maximum. C'est un plancher physique.

Trois façons de le contourner, et il n'y en a pas d'autres :

Levier	Effet	Chapitre
Quantifier	divise \(Q\) par 2 ou 4	IA · Quantification
Grouper les requêtes	amortit \(Q\) sur plusieurs jetons, \(I\) croît avec le lot	IA · Entraînement vs inférence
Réduire le nombre de poids lus	MoE : on ne lit que les experts actifs	IA · MoE

Et une quatrième, qui ne change pas \(I\) mais élimine le surcoût qui empêche d'atteindre le plancher : les megakernels. C'est exactement leur promesse — Hazy Research atteint 78 % de la bande passante contre ~50 % pour les systèmes existants.

1.4 Le roofline hiérarchique¶

Le modèle de base ne considère que la mémoire globale. Or il y a des plafonds à chaque niveau.

performance
    │              ┌───────────  pic tensor core
    │             ╱┌────────────  pic FP32
    │            ╱╱
    │       ╱───╱╱      ← pente L1/shared (~30 To/s)
    │      ╱  ╱╱
    │     ╱  ╱  ← pente L2 (~7 To/s)
    │    ╱  ╱
    │   ╱  ╱ ← pente HBM (3,35 To/s)
    └──┴──┴────────────────────→ I

Chaque niveau de mémoire a sa propre droite. Un noyau peut être limité par la bande passante de la mémoire partagée alors qu'il est loin du plafond HBM — c'est fréquent dans les GEMM naïves, où la boucle interne lit deux valeurs partagées par FMA.

Nsight Compute sait tracer ces rooflines hiérarchiques : la section roofline collecte les plafonds L1, L2 et mémoire globale sur le même graphique.

1.5 Les limites du modèle¶

Le roofline est un modèle, avec des hypothèses. Il faut savoir quand il ment.

Ce que le roofline ne capture pas

1. La latence. Le modèle suppose que la bande passante est saturée. Si vous n'avez pas assez de parallélisme, vous serez limité par la latence bien avant d'approcher \(B\). C'est le sujet du chapitre 3.

2. L'efficacité des transferts. \(Q\) est le volume utile. Un accès non coalescé transfère jusqu'à 8× plus d'octets que nécessaire. Le roofline calculé sur les octets utiles surestime alors la performance atteignable.

3. Le surcoût de lancement. Pour un noyau de 20 µs, 5 µs de lancement représentent 25 % du temps et n'apparaissent nulle part dans le modèle.

4. Les caches. Si vos données tiennent en L2, le \(B\) pertinent est celui du L2, pas de la HBM. Le modèle simple donne alors une borne trop pessimiste.

5. Les unités spécialisées. Un noyau qui sature la SFU (exponentielles) n'est ni memory-bound ni compute-bound au sens du modèle. C'est exactement le cas de l'attention sur Blackwell, où FlashAttention-4 déplace les exponentielles vers les unités FMA parce que la SFU sature.

Malgré tout cela, le roofline reste le premier outil à sortir. Il vous dit dans quelle direction chercher, ce qui est 80 % du travail.

1.6 Un formulaire à garder¶

Calculer \(Q\) — comptez les octets strictement nécessaires, en supposant une réutilisation parfaite :

Q = Σ (taille de chaque tenseur lu au moins une fois)
  + Σ (taille de chaque tenseur écrit)

Calculer \(W\) — comptez les opérations flottantes :

Opération	FLOP
`a + b`, `a * b`	1
`fma(a,b,c)`	2
GEMM \(M{\times}K \cdot K{\times}N\)	\(2MNK\)
`exp`, `log`, `rsqrt`	1 (mais sur la SFU, ~4× plus lent)
Softmax sur \(n\) éléments	~5n
LayerNorm sur \(n\) éléments	~7n
Attention, séq. \(S\), dim. \(d\)	\(4S^2 d\) (les deux GEMM)

Comparer :

\[ \text{ratio} = \frac{I}{I_{\text{crit}}} \]

ratio \(< 0{,}1\) → franchement limité par la mémoire, ne touchez pas au calcul ;
\(0{,}1 <\) ratio \(< 10\) → zone mixte, les deux leviers comptent ;
ratio \(> 10\) → limité par le calcul, utilisez les tensor cores.

Résumé du chapitre¶

À retenir

\(I = W/Q\), \(I_{\text{crit}} = P/B\). Comparer les deux donne le régime.
\(I_{\text{crit}}\) vaut ~296 sur H100 en BF16 et monte à chaque génération : de plus en plus de noyaux sont limités par la mémoire.
L'intensité arithmétique est une propriété de l'implémentation, pas du problème. Le pavage en tuiles la fait croître.
Le décodage LLM à lot 1 a \(I \approx 1\) : plancher physique de ~42 ms par jeton pour un 70B en BF16 sur H100.
Mesurez en Go/s si limité par la mémoire, en TFLOPS si limité par le calcul. Jamais l'inverse.
Le modèle ignore la latence, l'efficacité des transferts, le surcoût de lancement et les unités spécialisées. Il indique la direction, pas la destination.

Vérifiez que vous avez compris¶

Une convolution 3×3 sur une image de 1024×1024 en float. Limitée par quoi sur H100 ?

\(W = 1024^2 \times 9 \times 2 = 1{,}89 \times 10^7\) FLOP ;
\(Q\) : si l'on réutilise parfaitement (chaque pixel lu une fois grâce à la mémoire partagée), \(Q = 2 \times 1024^2 \times 4 = 8{,}39 \times 10^6\) octets ;
\(I = 2{,}25\).

Contre \(I_{\text{crit}} = 20\) en FP32 : limitée par la mémoire, d'un facteur ~9.

Sans réutilisation (implémentation naïve), chaque pixel est lu 9 fois : \(Q = 3{,}8\times 10^7\), \(I = 0{,}5\), soit 40× sous le seuil. D'où l'importance de la mémoire partagée ou de la mémoire de texture ici.

Votre GEMM atteint 400 TFLOPS en BF16 sur H100. Bon ou mauvais ?

400 / 990 = 40 % du pic. C'est correct pour une implémentation manuelle, médiocre par rapport à cuBLAS ou CUTLASS qui dépassent 80 % sur de grandes matrices.

Les causes probables, dans l'ordre : pas de wgmma (l'API wmma plafonne à ~60 %), pas de TMA donc pipelining insuffisant, quantification de vagues, tuiles trop petites. Voir IA · GEMM.

Pourquoi le décodage LLM devient-il limité par le calcul quand le lot grandit ?

Parce que les poids sont lus une seule fois pour tout le lot. Avec un lot de taille \(b\) :

\[Q \approx 2\Theta \quad\text{(inchangé)}, \qquad W \approx 2\Theta b\]

donc \(I \approx b\). Il faut \(b \approx 296\) pour atteindre le seuil sur H100 en BF16.

C'est la raison d'être du continuous batching de vLLM et SGLang : sans lot, on gaspille 99,7 % des tensor cores. C'est aussi pourquoi les megakernels latence (lot 1) et les megakernels débit (grand lot) sont des objets techniques différents — voir Megakernels 3.

Chapitre suivant : 2 · Coalescence et conflits de banc

Sources de ce chapitre¶

Williams, Waterman, Patterson, Roofline: An Insightful Visual Performance Model for Multicore Architectures, CACM 2009 — DOI 10.1145/1498765.1498785.
Modal GPU Glossary — Roofline model
Hierarchical Roofline Performance Analysis for Deep Learning Applications, arXiv:2009.05257
Nsight Compute Profiling Guide — section Roofline
Hazy Research, Look Ma, No Bubbles! — 78 % de bande passante contre ~50 % pour les systèmes existants.