Aller au contenu

5 · Plan d'entraînement de 90 jours

Douze semaines, avec un livrable vérifiable par semaine. Conçu pour environ une heure par jour en semaine, plus quelques heures le week-end.

Le principe

Trois règles.

1. Un livrable par semaine. Pas « avoir lu », mais « avoir écrit un programme qui fait X et l'avoir mesuré ».

2. Mesurer avant d'optimiser. Chaque livrable inclut un chiffre : Go/s, TFLOPS, ou fraction du plafond.

3. Ne pas sauter les semaines 1 à 4. Elles sont ennuyeuses et tout le reste en dépend.

Phase 1 — Les fondations (semaines 1-4)

Semaine 1 : le premier noyau

Lire : Fondations 1-3, Modèle 1-2

Écrire :

  • SAXPY complet, avec vérification d'erreurs ;
  • version avec boucle à parcours de grille ;
  • mesure du temps avec des événements CUDA.

Livrable : « Mon SAXPY atteint X Go/s sur Y possibles, soit Z %. »

Réussite si : Z > 80 %.

Semaine 2 : la mémoire

Lire : Fondations 4, Performance 1-2

Écrire :

  • transposition de matrice, version naïve ;
  • version avec mémoire partagée ;
  • version avec padding pour éliminer les conflits de banc.

Livrable : un tableau des trois versions avec les Go/s de chacune.

Réussite si : la progression est ~1× → 4× → 5×, et vous savez expliquer chaque saut.

L'exercice le plus formateur du plan

La transposition est le meilleur exercice pour comprendre la mémoire : elle n'a aucun calcul, donc tout ce que vous mesurez est du mouvement de données.

Semaine 3 : les réductions

Lire : Modèle 3-5

Écrire : les six versions de la réduction, de la naïve à celle à deux niveaux avec shuffles.

Livrable : un graphique des six versions.

Réussite si : la dernière version dépasse 85 % de la bande passante.

Semaine 4 : le profilage

Lire : Performance 3-5

Faire :

  • profiler les noyaux des semaines 1 à 3 avec Nsight Compute ;
  • pour chacun, relever Speed of Light, les états de warp et l'occupancy ;
  • rédiger une explication de ce qui limite chaque noyau.

Livrable : trois rapports d'une page.

Réussite si : vos explications correspondent aux mesures.

Phase 2 — Le calcul (semaines 5-8)

Semaine 5 : la GEMM par tuiles

Lire : IA 2, Modèle 3

Écrire : GEMM naïve, puis coalescée, puis par tuiles en mémoire partagée.

Livrable : trois versions, mesurées en TFLOPS, comparées à cuBLAS.

Réussite si : la version par tuiles atteint 25 à 35 % de cuBLAS en FP32.

Semaine 6 : la GEMM optimisée

Écrire : pavage en registres 1D puis 2D, chargements vectorisés.

Livrable : « Ma GEMM atteint X % de cuBLAS. »

Réussite si : X > 70 %.

Le repère

Simon Boehm atteint 95 % de cuBLAS en dix étapes. Si vous êtes à 70 % en deux semaines, c'est un très bon résultat.

Semaine 7 : Triton

Lire : Écosystème 1-2

Écrire :

  • addition vectorielle en Triton ;
  • softmax fusionné ;
  • RMSNorm fusionné ;
  • GEMM en Triton, avec autotuning.

Livrable : comparaison Triton contre vos noyaux CUDA de la semaine 6.

Réussite si : la GEMM Triton dépasse votre GEMM CUDA, et vous comprenez pourquoi.

Semaine 8 : l'intégration

Lire : Pratique 3-4

Écrire :

  • un opérateur custom enregistré via torch.library ;
  • son gradient, vérifié par gradcheck ;
  • une suite de tests couvrant les cinq cas ;
  • une mesure de bout en bout dans un petit modèle.

Livrable : un dépôt installable avec pip install -e . et des tests qui passent.

Réussite si : torch.compile traverse votre opérateur sans graph break.

Phase 3 — L'IA (semaines 9-10)

Semaine 9 : l'attention

Lire : IA 1, 3

Écrire : FlashAttention simplifiée en Triton — pavage, softmax en ligne, masque causal.

Livrable : comparaison contre torch.nn.functional.scaled_dot_product_attention, à plusieurs longueurs de séquence.

Réussite si : correct à \(10^{-2}\) près en BF16, et l'écart de performance avec la référence reste inférieur à 3×.

Ne visez pas la parité avec FlashAttention

FlashAttention est le produit de quatre années de travail par des spécialistes. Atteindre le tiers de sa performance en une semaine est déjà excellent.

L'objectif est de comprendre l'algorithme, pas de le battre.

Semaine 10 : la quantification et la fusion

Lire : IA 4-5

Écrire :

  • un noyau de déquantification INT4 → BF16 ;
  • une GEMM W4A16 simple ;
  • une fusion RMSNorm + quantification.

Livrable : mesure du gain en bande passante par rapport à la version BF16.

Réussite si : le trafic HBM mesuré (dram__bytes.sum) est bien divisé par ~4.

Phase 4 — Les megakernels (semaines 11-12)

Semaine 11 : le noyau persistant

Lire : Megakernels 1-3, CUDA moderne 1

Écrire :

  • un noyau persistant avec file de travail dynamique (atomicAdd) ;
  • une barrière de grille manuelle avec compteur de génération ;
  • deux opérations chaînées dans un seul noyau, avec un compteur de dépendance.

Livrable : « Deux opérations en un lancement, résultat correct, X µs économisées par rapport à deux lancements. »

Réussite si : correct, et compute-sanitizer --tool racecheck est propre.

Semaine 12 : le mini-megakernel

Lire : Megakernels 8-9

Écrire : un megakernel exécutant une couche de transformeur complète — RMSNorm, QKV, attention, projection O, MLP — en un seul lancement, avec allocateur de pages et compteurs de dépendance.

Livrable : comparaison contre la version à noyaux multiples, en latence.

Réussite si : correct jeton à jeton, et vous savez expliquer d'où vient (ou ne vient pas) le gain.

Attente réaliste pour la semaine 12

Votre mini-megakernel sera probablement plus lent que la version à noyaux multiples, parce que vos instructions individuelles seront moins optimisées que cuBLAS et FlashAttention.

Ce n'est pas un échec. L'objectif est de comprendre les mécanismes : persistance, dispatch, pages, compteurs. Le gain vient ensuite, quand les instructions atteignent le niveau des noyaux qu'elles remplacent — et cela demande des mois.

Le tableau de suivi

Sem. Sujet Livrable Critère
1 Premier noyau SAXPY > 80 % de la bande passante
2 Mémoire Transposition ×3 progression 1× → 4× → 5× expliquée
3 Réductions Réduction ×6 > 85 % de la bande passante
4 Profilage 3 rapports explications = mesures
5 GEMM tuiles GEMM ×3 25-35 % de cuBLAS
6 GEMM registres GEMM optimisée > 70 % de cuBLAS
7 Triton 4 noyaux Triton > CUDA manuel
8 Intégration Extension PyTorch pas de graph break
9 Attention FlashAttention Triton correct, < 3× de l'écart
10 Quantification GEMM W4A16 trafic HBM ÷ 4
11 Persistant 2 ops fusionnées racecheck propre
12 Megakernel 1 couche complète correct jeton à jeton

Après les 90 jours

Trois directions, selon ce qui vous a le plus intéressé.

Approfondir les noyaux

  • Lire CUTLASS et les tutoriels Colfax.
  • Refaire la GEMM avec TMA et wgmma sur un H100 loué.
  • Participer aux leaderboards GPU MODE.

Approfondir les systèmes

  • Lire le code de vLLM ou SGLang.
  • Essayer Mirage MPK sur un modèle.
  • Reproduire une partie des résultats de la partie 8.

Approfondir les compilateurs

  • Lire le code de Triton, ses passes MLIR.
  • Écrire un backend Triton ou une passe d'optimisation.
  • Suivre les travaux sur CUDA Tile et TIRx.

Les erreurs à éviter

Les cinq façons d'échouer

1. Tout lire avant d'écrire. 34 heures de lecture sans code = zéro rétention. Alternez.

2. Ne pas mesurer. Un noyau non mesuré est une hypothèse. Chaque livrable a un chiffre.

3. Optimiser trop tôt. D'abord correct, puis mesuré, puis optimisé. Dans cet ordre.

4. Comparer à la mauvaise référence. « 10× plus rapide que NumPy » ne veut rien dire. Comparez à cuBLAS, cuDNN, PyTorch.

5. Sauter le profilage. La semaine 4 est celle qu'on veut sauter et celle qui change tout. Sans elle, vous optimisez à l'aveugle pendant huit semaines.

Résumé de la partie 9

À retenir

  • Colab gratuit couvre 70 % du plan ; louez quelques heures de H100 pour les semaines 11-12.
  • Triton est le chemin par défaut pour intégrer un noyau ; torch.library.custom_op est ce qui évite les graph breaks.
  • Tester : cinq cas, tolérances adaptées à la précision, jamais torch.equal, compute-sanitizer systématique.
  • Douze semaines, douze livrables mesurés. Un chiffre par semaine.
  • Attente réaliste : votre mini-megakernel sera plus lent que la version à noyaux multiples. C'est normal et ce n'est pas le point.

Partie suivante : Ressources

Sources de ce chapitre