Développement GPU : de zéro à expert¶
Cours complet de programmation sur processeur graphique, depuis les principes physiques qui donnent au GPU sa forme jusqu'aux techniques d'écriture de noyaux qui définissent l'état de l'art en 2026.
Ce que ce document est¶
Un cours progressif, pas une documentation de référence. Chaque notion est introduite avant d'être utilisée, chaque symbole est défini, chaque chiffre est sourcé. Il se lit dans l'ordre, mais chaque partie possède son propre index et peut servir de point d'entrée à un lecteur déjà formé.
Il couvre :
- le matériel — ce qu'est un SM, un warp, un tensor core, pourquoi la hiérarchie mémoire a la forme qu'elle a, et quels chiffres il faut connaître par cœur ;
- le modèle de programmation — CUDA d'abord, parce que c'est le modèle que tous les autres imitent ;
- la performance — le modèle roofline, la coalescence, l'occupancy, le profilage, et une méthode de travail reproductible ;
- le CUDA moderne — asynchronisme, TMA, clusters,
wgmma,tcgen05, spécialisation des warps, CUDA Tile ; - l'écosystème — Triton, Gluon, CUTLASS/CuTe DSL, ThunderKittens, Helion, Mojo, HIP/ROCm, SYCL, WebGPU, Metal ;
- tous les domaines d'application — graphique temps réel, calcul scientifique, bases de données, signal et image, finance, cryptographie ;
- l'intelligence artificielle en profondeur — GEMM, FlashAttention, fusion, quantification, MoE, collectives multi-GPU ;
- les megakernels — dix chapitres sur l'idée qui structure la recherche en systèmes d'inférence depuis 2025 ;
- la pratique — obtenir une carte, monter un environnement, intégrer un noyau dans PyTorch, tester, déboguer, et un plan d'apprentissage de 90 jours ;
- les ressources — livres, cours vidéo, papiers, blogs et dépôts, classés et commentés.
Ce que ce document n'est pas¶
Ce n'est pas un remplacement du CUDA C++ Programming Guide, ni du PTX ISA. Ces documents font plusieurs milliers de pages et sont la référence normative. Ce cours vous apprend à les lire et vous dit où regarder.
Ce n'est pas non plus un catalogue de recettes. Les recettes changent à chaque génération de matériel ; les modèles mentaux, non.
Public visé¶
| Profil | Ce que vous y gagnez |
|---|---|
| Développeur qui n'a jamais écrit de code GPU | Le parcours complet, dans l'ordre |
| Utilisateur de PyTorch qui veut écrire ses noyaux | Parties 2, 3, 5 et 7 |
| Ingénieur performance déjà formé au CUDA | Parties 4, 7 et 8 |
| Chercheur en systèmes d'apprentissage | Partie 8 (megakernels), puis les sources |
| Étudiant | Le parcours complet + le plan 90 jours (partie 10) |
Prérequis réels : savoir programmer dans un langage impératif (Python et un peu de C suffisent), et connaître la notion de tableau en mémoire. Tout le reste est construit dans le document.
Ce qui n'est pas prérequis : l'algèbre linéaire avancée, l'architecture des processeurs, l'apprentissage automatique. Les notions nécessaires sont introduites au moment où elles servent.
Temps de lecture et difficulté¶
| Partie | Chapitres | Lecture | Difficulté |
|---|---|---|---|
| Fondations matérielles | 6 | 3 h | ★☆☆ |
| Modèle de programmation | 7 | 4 h | ★★☆ |
| Performance | 5 | 3 h | ★★☆ |
| CUDA moderne | 6 | 4 h | ★★★ |
| Écosystème | 9 | 4 h | ★★☆ |
| Domaines d'application | 6 | 2 h 30 | ★☆☆ |
| GPU pour l'IA | 8 | 5 h | ★★★ |
| Megakernels | 10 | 6 h | ★★★ |
| Pratique | 5 | 2 h | ★★☆ |
| Ressources | 4 | 1 h | ★☆☆ |
| Annexes | 3 | — | — |
Total : environ 34 h de lecture attentive. Le document n'est pas conçu pour être lu d'une traite.
Parcours de lecture¶
Voir le parcours de lecture détaillé, qui propose six trajectoires selon votre point de départ et votre objectif.
Si vous êtes pressé : lisez le résumé exécutif, puis la partie Megakernels.
Les onze parties¶
1 · Fondations matérielles¶
Pourquoi un GPU existe, comment il est construit, et les six chiffres qu'il faut connaître. On y construit le modèle mental « débit contre latence » qui explique tout le reste.
2 · Modèle de programmation¶
De hello world à la réduction parallèle. Grille, blocs, warps, mémoire
partagée, primitives de warp, atomiques, flux et graphes.
3 · Performance¶
Le modèle roofline, la coalescence, les conflits de banc, l'occupancy et son mythe, le profilage avec Nsight, et une checklist opérationnelle.
4 · CUDA moderne¶
Ce qui a changé depuis Ampere : copies asynchrones, mbarrier, Tensor Memory
Accelerator, clusters et mémoire partagée distribuée, wgmma puis tcgen05,
spécialisation des warps, et le modèle CUDA Tile de 2025-2026.
5 · Écosystème¶
Neuf façons d'écrire un noyau en 2026, et comment choisir : CUDA C++, Triton, Gluon, CUTLASS et CuTe DSL, ThunderKittens, Helion, Mojo, HIP, et la famille portable (SYCL, OpenCL, Vulkan, WebGPU, Metal).
6 · Domaines d'application¶
Le GPU ailleurs que dans l'IA : rendu temps réel, calcul scientifique, traitement de données, signal et image, finance, cryptographie — et une liste honnête de ce qui ne va pas sur GPU.
7 · GPU pour l'IA¶
Le cœur applicatif. Anatomie d'un transformeur du point de vue du GPU, GEMM optimisée pas à pas, FlashAttention, fusion de noyaux, quantification, MoE, collectives multi-GPU, et la différence structurelle entre entraînement et inférence.
8 · Megakernels¶
Le focus. Pourquoi les frontières de noyau coûtent cher, comment on les supprime, et l'état de l'art complet : l'interpréteur sur GPU de Stanford, Mirage MPK, Event Tensor, Fleet, le monokernel AMD, et comment en écrire un.
9 · Pratique¶
Obtenir un GPU sans en acheter un, monter l'environnement, intégrer un noyau dans PyTorch, tester la correction, déboguer, et un plan d'apprentissage de 90 jours.
10 · Ressources¶
Livres, cours vidéo, papiers de recherche, blogs et dépôts — classés, commentés, avec ce qu'il faut en attendre.
11 · Annexes¶
Glossaire, tableau comparatif du matériel, et la liste complète des sources.
Avertissements méthodologiques¶
Le matériel change plus vite que les modèles mentaux
Les chiffres de ce document sont ceux d'août 2026 : Hopper (H100/H200) et
Blackwell (B200/GB200) sont les plateformes de référence, Rubin entre en
production. Une instruction PTX comme tcgen05.mma n'existe que sur
sm_100. Les principes — latence cachée par le parallélisme, coût des
accès non coalescés, arithmétique du roofline — sont stables depuis 2007.
Les chiffres de performance publiés sont des revendications
Quand ce document rapporte « 1,7× plus rapide que vLLM », c'est ce que le papier annonce, dans sa configuration, avec sa version de la ligne de base. Les accélérations de noyaux sont particulièrement sensibles au choix de la référence. Chaque chiffre est attribué à sa source ; aucun n'a été reproduit indépendamment ici.
Le code de ce document n'a pas été exécuté sur GPU
L'environnement de rédaction ne dispose ni de GPU ni de compilateur CUDA. Les extraits de code sont écrits pour être lus et sont conformes aux documentations citées, mais ils n'ont pas été compilés ni mesurés. Ils sont signalés comme tels quand le risque d'erreur est réel. Les seuls chiffres calculés dans ce document sont arithmétiques (arithmétique d'intensité, budgets mémoire) et refaisables au papier.
Créé le 22 août 2026. Dernière révision : 22 août 2026.
Principales sources : CUDA C++ Programming Guide · Programming Massively Parallel Processors, 5e éd. (2026) · Mirage Persistent Kernel (arXiv:2512.22219) · Hazy Research, Look Ma, No Bubbles! · GPU MODE · liste complète en annexes/sources.