Aller au contenu

5 · Checklist d'optimisation

La liste opérationnelle. Ordonnée par rendement décroissant, avec pour chaque point le symptôme qui le déclenche et le gain typique.

Comment utiliser cette liste

Ne la parcourez pas dans l'ordre. Profilez d'abord, puis sautez au point que le symptôme désigne. Les points sont numérotés par rendement moyen, pas par ordre d'application.

Après chaque changement : mesurez. Une optimisation non mesurée est une hypothèse.

Niveau 0 — Avant d'optimiser quoi que ce soit

# Action Pourquoi
0.1 Vérifier que le résultat est correct Un noyau faux et rapide n'a aucune valeur. Écrire le test avant.
0.2 Établir une ligne de base mesurée Sans référence, « plus rapide » n'a pas de sens.
0.3 Calculer le roofline au papier Deux minutes qui vous disent où chercher.
0.4 Vérifier que le GPU travaille (Nsight Systems) Optimiser un noyau pendant que le GPU est inactif 70 % du temps est du gaspillage.
0.5 Vérifier les flags de compilation -O3, -arch correct, -lineinfo, jamais -G pour mesurer.

Le point 0.1 n'est pas une formalité

La grande majorité des optimisations GPU introduisent des bugs de synchronisation ou de bord. Un test qui compare à une référence CPU, exécuté après chaque changement, est le seul filet.

torch.testing.assert_close(mon_noyau(x), reference(x), rtol=1e-2, atol=1e-2)

Niveau 1 — Réduire le travail (gain : 2× à 100×)

Le meilleur code est celui qui ne s'exécute pas.

# Action Symptôme déclencheur Gain typique
1.1 Fusionner les noyaux élémentaires plusieurs petits noyaux consécutifs sur les mêmes données 2× à 5×
1.2 Réduire la précision limité par la mémoire, précision non critique 1,5× à 4×
1.3 Éliminer les tenseurs intermédiaires trafic HBM >> volume utile 2× à 10×
1.4 Choisir un meilleur algorithme complexité asymptotique améliorable illimité
1.5 Ne pas calculer ce qui est masqué attention causale, padding jusqu'à 2×

1.1 — La fusion, en pratique

# 3 noyaux, 3 allers-retours en HBM
y = torch.relu(x)
y = y * scale
y = y + biais

# 1 noyau
y = torch.compile(lambda x: torch.relu(x) * scale + biais)(x)

Pour trois opérations élémentaires sur un tenseur de \(n\) éléments en float :

  • non fusionné : \(3 \times (4n + 4n) = 24n\) octets ;
  • fusionné : \(4n + 4n = 8n\) octets.

Facteur 3 exactement, et il croît linéairement avec le nombre d'opérations fusionnées.

C'est la version « locale » de l'idée qui, poussée jusqu'à son terme sur un modèle entier, donne les megakernels.

1.5 — L'exemple du masque causal

Dans une attention causale, la moitié supérieure de la matrice \(\mathbf{Q}\mathbf{K}^\top\) est masquée. Une implémentation naïve calcule les \(S^2\) produits puis en jette la moitié. FlashAttention saute directement les blocs entièrement masqués : facteur 2 gratuit.

Niveau 2 — Améliorer les accès mémoire (gain : 1,5× à 8×)

# Action Symptôme Gain typique
2.1 Coalescer secteurs/requête > 4 jusqu'à 8×
2.2 Passer en SoA structures avec plusieurs champs 2× à 6×
2.3 Vectoriser (float4) limité par la mémoire, accès contigus 1,1× à 1,3×
2.4 Utiliser la mémoire partagée données relues plusieurs fois 2× à 30×
2.5 Éliminer les conflits de banc bank_conflicts non nul jusqu'à 32× sur l'accès
2.6 Aligner les allocations secteurs/requête = 5 au lieu de 4 1,25×
2.7 Trier les indirections accès a[idx[i]] 2× à 5×
2.8 Marquer le non-temporel données lues une seule fois 1,1× à 1,3×

2.4 — Le seuil de rentabilité

La mémoire partagée n'est rentable que si une donnée est relue au moins deux fois. Sinon, on paie un aller-retour supplémentaire pour rien.

Règle : si chaque élément chargé est utilisé \(r\) fois, le gain est d'environ \(\min(r, B_{\text{shared}}/B_{\text{HBM}}) \approx \min(r, 9)\) sur H100.

Niveau 3 — Améliorer l'exécution (gain : 1,2× à 3×)

# Action Symptôme Gain typique
3.1 Utiliser les tensor cores GEMM ou convolution en FP32 8× à 30×
3.2 Augmenter l'ILP Stall Long Scoreboard avec occupancy élevée 1,3× à 2×
3.3 Dérouler les boucles boucle courte à borne connue 1,1× à 1,5×
3.4 Réduire la divergence Branch Efficiency < 90 % jusqu'à 2×
3.5 Agréger les atomiques forte contention 5× à 50× sur l'opération
3.6 Ajuster la taille de bloc balayage 1,05× à 1,3×
3.7 Éviter le spilling Local Memory non nul dans une boucle 1,5× à 5×
3.8 Corriger la quantification de vagues nombre de blocs ≈ 1,1 × blocs par vague jusqu'à 1,9×
3.9 Utiliser les intrinsèques rapides expf, sinf, division 1,1× à 3× sur l'opération

3.9 — Les intrinsèques

Standard Rapide Précision
expf(x) __expf(x) ~2 ULP
logf(x) __logf(x) ~2 ULP
sinf(x) __sinf(x) ~2 ULP
a / b __fdividef(a,b) ~2 ULP
sqrtf(x) __fsqrt_rn(x) exacte
1/sqrtf(x) rsqrtf(x) ~2 ULP, une instruction

Ou globalement : nvcc -use_fast_math, qui active aussi -ffast-math, -prec-div=false et le flush-to-zero des dénormaux. Puissant et à manier avec précaution en calcul scientifique.

Niveau 4 — CUDA moderne (gain : 1,3× à 3×, sur Hopper/Blackwell)

# Action Prérequis Gain typique
4.1 Copies asynchrones (cp.async) Ampere+ 1,2× à 1,5×
4.2 TMA Hopper+, sm_90a 1,2× à 2×
4.3 Spécialisation des warps Hopper+ 1,3× à 2×
4.4 Clusters + DSMEM Hopper+ 1,1× à 1,4×
4.5 wgmma / tcgen05 via CUTLASS ou CuTe DSL 1,5× (vs mma)
4.6 CUDA Graphs tous 1,1× à 1,4× sur des noyaux courts
4.7 Persistance L2 Ampere+ 1,1× à 1,3×

Tous ces points sont traités en partie 4.

Le rapport effort/gain se dégrade vite

Les points de niveau 4 sont beaucoup plus difficiles à mettre en œuvre que ceux de niveau 1 et 2, pour des gains bien inférieurs. Ne les abordez qu'après avoir épuisé les niveaux précédents, et de préférence via une bibliothèque (CUTLASS, Triton, ThunderKittens) plutôt qu'à la main.

Niveau 5 — Structurel (gain : variable, effort important)

# Action Quand Référence
5.1 Noyau persistant beaucoup de petits noyaux, lot faible Megakernels 2
5.2 Megakernel inférence LLM latence-critique Megakernels
5.3 Recouvrir calcul et communication multi-GPU IA 7
5.4 Localité de chiplet B200, MI300X/MI350X Megakernels 6
5.5 Changer de bibliothèque avant tout le reste Écosystème

Le point 5.5 devrait être le premier

Avant d'écrire quoi que ce soit, vérifiez que cuBLAS, cuDNN, CUTLASS, FlashAttention, ou une bibliothèque du domaine ne fait pas déjà le travail. Vous ne battrez pas cuBLAS sur une GEMM générique, et ce n'est pas une faiblesse : c'est le résultat de vingt ans-personnes d'optimisation.

La bonne question n'est pas « comment écrire une GEMM rapide » mais « quelle opération de mon programme n'a pas de bibliothèque ? ». C'est presque toujours une fusion ou une forme inhabituelle.

La checklist express

Pour un noyau donné, dix questions :

□  Ai-je un test de correction qui passe ?
□  Ai-je une mesure de référence ?
□  Quelle est l'intensité arithmétique ? Limité par quoi ?
□  Le GPU est-il occupé (Nsight Systems) ?
□  Speed of Light : compute % et memory % ?
□  Secteurs par requête = 4 ?
□  Conflits de banc = 0 ?
□  Spilling = 0 ?
□  Les tensor cores sont-ils utilisés (chercher HMMA dans le SASS) ?
□  Le nombre de blocs est-il un multiple des blocs par vague ?

Si les dix cases sont cochées et que vous êtes à plus de 80 % du plafond pertinent, arrêtez. Le temps restant est mieux investi ailleurs.

Les cinq pièges qui invalident une mesure

À vérifier avant de croire un chiffre

  1. Pas d'échauffement — la première exécution inclut le JIT et une fréquence basse.
  2. Données en cache — un noyau qui relit 20 Mo depuis le L2 semble atteindre 10 To/s.
  3. Fréquences non verrouillées — la variabilité thermique donne ±15 %.
  4. Ligne de base malhonnête — comparer à une version non optimisée de soi-même, ou à NumPy.
  5. Pas de synchronisation — mesurer la mise en file au lieu de l'exécution.

Ces cinq points expliquent une grande partie des accélérations spectaculaires annoncées dans la littérature et sur les réseaux.

Résumé de la partie 3

Les cinq idées à emporter

  1. Calculez le roofline avant de coder. \(I = W/Q\) contre \(I_{\text{crit}} = P/B\).
  2. Le plus gros gain est de ne pas transférer : fusion, précision réduite, mémoire partagée.
  3. La coalescence vaut jusqu'à 8×, les conflits de banc jusqu'à 32×, et les deux sont invisibles dans le résultat.
  4. L'occupancy est un moyen, pas un but. Diagnostiquez avec les états de warp.
  5. Speed of Light dans Nsight Compute suffit dans 80 % des cas.

Partie suivante : CUDA moderne

Sources de ce chapitre