Aller au contenu

Parcours de lecture

Ce document fait 69 chapitres. Personne ne devrait les lire dans l'ordre sans savoir pourquoi. Cette page propose six trajectoires.

Comment choisir

Répondez à deux questions.

1. Qu'est-ce que vous savez déjà ?

Vous savez… Point de départ
Programmer, mais rien du GPU Fondations, chapitre 1
Utiliser PyTorch, jamais écrit de noyau Fondations, chapitre 2
Écrire un noyau CUDA simple Performance
Optimiser un noyau, profiler CUDA moderne
Écrire des noyaux Hopper avec TMA Megakernels

2. Qu'est-ce que vous voulez en faire ?

Objectif Parcours
Comprendre, sans écrire de code A
Écrire mon premier noyau utile B
Rendre PyTorch plus rapide C
Faire du calcul scientifique D
Travailler sur l'inférence LLM E
Comprendre les megakernels précisément F

Parcours A — Comprendre sans coder

≈ 5 h. Pour un lecteur qui veut le modèle mental sans la pratique : décideur technique, chercheur d'un autre domaine, curieux.

  1. Résumé exécutif
  2. Fondations — chapitres 1 à 5
  3. Performance · Le modèle roofline
  4. Écosystème · Panorama des langages
  5. Domaines d'application — la partie entière
  6. Megakernels — chapitres 1, 2 et 10
  7. Annexes · Glossaire

Vous saurez expliquer pourquoi un GPU est rapide, pourquoi il est parfois lent, et de quoi parlent les gens qui écrivent des noyaux.

Parcours B — Premier noyau utile

≈ 12 h, plus le temps de coder. Le parcours canonique du débutant.

  1. Fondations — tout, dans l'ordre
  2. Modèle de programmation — tout
  3. Pratique · Obtenir un GPU
  4. Pratique · Environnement
  5. Performance — tout
  6. Pratique · Tester et déboguer
  7. Pratique · Plan 90 jours

À la fin, vous devez avoir écrit une réduction, une transposition et une multiplication matricielle par blocs, et les avoir profilées.

Le test de sortie du parcours B

Écrire une transposition de matrice qui atteint plus de 80 % de la bande passante mémoire de votre carte, et savoir expliquer, chiffres du profileur à l'appui, pourquoi elle n'atteint pas 100 %.

Parcours C — Accélérer PyTorch

≈ 10 h. Pour l'ingénieur ML qui veut sortir du cadre des opérateurs existants.

  1. Fondations — chapitres 1 à 4
  2. Modèle de programmation — chapitres 1 à 3
  3. Performance · Roofline et Profilage
  4. Écosystème · Triton et Gluon
  5. Écosystème · Helion et torch.compile
  6. IA · Fusion de noyaux
  7. Pratique · Intégrer un noyau dans PyTorch
  8. IA · Quantification si c'est votre sujet

Triton avant CUDA est un choix assumé ici : pour la fusion élémentaire, le rapport effort/gain est bien meilleur, et le modèle de tuiles prépare directement au CUDA moderne.

Parcours D — Calcul scientifique

≈ 12 h. Pour qui vient du HPC, de la simulation ou du traitement de données.

  1. Fondations — tout
  2. Modèle de programmation — tout
  3. Performance — tout
  4. Fondations · Nombres flottants, à relire attentivement
  5. Domaines · Calcul scientifique
  6. Domaines · Ce qui ne va pas sur GPU
  7. Écosystème · Portabilité et AMD/ROCm
  8. IA · Multi-GPU et collectives

Le chapitre sur les flottants est ici obligatoire : c'est le domaine où la précision compte réellement, et où les raccourcis de l'IA ne sont pas transposables.

Parcours E — Inférence LLM

≈ 15 h. Le parcours le plus demandé en 2026.

  1. Fondations — chapitres 1, 2, 4, 5, 6
  2. Modèle de programmation — chapitres 1 à 4
  3. Performance · Roofline
  4. CUDA moderne — tout
  5. IA — tout, c'est le cœur
  6. Megakernels — tout
  7. Ressources · Papiers

Parcours F — Megakernels, précisément

≈ 8 h. Pour un lecteur qui maîtrise déjà CUDA et veut la partie 8 sans préliminaires.

Prérequis vérifiés : vous savez ce qu'est un warp, une mbarrier, une copie TMA, et pourquoi cudaDeviceSynchronize() entre deux noyaux est une mauvaise idée. Si l'un des quatre manque, lisez d'abord CUDA moderne.

  1. Megakernels · Le problème des frontières de noyau
  2. Megakernels · Définition et généalogie
  3. Megakernels · L'interpréteur sur GPU
  4. Megakernels · Mirage MPK
  5. Megakernels · Event Tensor
  6. Megakernels · Fleet et les chiplets
  7. Megakernels · AMD
  8. Megakernels · Écrire le sien
  9. Megakernels · Quand ne pas le faire
  10. Megakernels · État de l'art

Notation utilisée dans tout le document

Symbole Signification
SM Streaming Multiprocessor, l'unité de calcul du GPU
CU Compute Unit, l'équivalent AMD
\(B\) bande passante mémoire, en octets par seconde
\(P\) débit de calcul crête, en opérations par seconde
\(I\) intensité arithmétique, en opérations par octet
\(t\) temps, en secondes (les µs et ns sont annotées)
sm_90, sm_100 cibles de compilation : Hopper, Blackwell
gfx942, gfx950 cibles AMD : CDNA 3 (MI300X), CDNA 4 (MI350X)

Les blocs de code sont annotés de leur langage. Un bloc marqué cuda est du CUDA C++ ; un bloc ptx est de l'assembleur virtuel NVIDIA.

Convention sur les mots « thread » et « fil »

Ce document utilise thread sans le traduire. « Fil d'exécution » est correct mais personne ne l'emploie dans ce domaine, et toute la littérature que vous lirez ensuite dit thread. De même pour warp, kernel (traduit ici par noyau quand le contexte est clair), tile (tuile) et occupancy (laissé tel quel : les traductions françaises sont ambiguës).

Après ce document

Trois directions, détaillées dans Ressources :

  • approfondir le matériel — les papiers de microbenchmarking (Volta, Turing, Hopper, Blackwell) démontent les architectures instruction par instruction ;
  • approfondir les compilateurs — MLIR, Triton, TVM, et la question ouverte de la génération automatique de noyaux ;
  • pratiquer — les puzzles GPU, KernelBench, les leaderboards GPU MODE.