Aller au contenu

Partie 4 · CUDA moderne

Ce qui a changé depuis Ampere, et pourquoi un noyau écrit selon les principes de 2018 laisse aujourd'hui la moitié de la machine inutilisée.

Le fil directeur : tout devient asynchrone

Le modèle CUDA classique est synchrone du point de vue du thread : une instruction de chargement bloque le thread jusqu'à ce que la donnée arrive. Le GPU cache cette latence en commutant vers d'autres warps.

Le modèle moderne inverse la charge : le thread déclenche une opération et continue. Il attend ensuite explicitement, au moment où il a réellement besoin du résultat.

Modèle classique                    Modèle moderne
────────────────                    ──────────────
ld  R1, [adr]     ← bloque          cp.async  [smem], [gmem]  ← ne bloque pas
                                    …autre travail…
use R1                              mbarrier.wait
                                    use smem

Le gain est double : moins de warps sont nécessaires pour saturer la mémoire, et le programmeur contrôle exactement ce qui est en vol.

Cette partie parcourt les mécanismes qui rendent cela possible, dans l'ordre où ils ont été introduits.

La chronologie

Génération Mécanisme Ce qu'il rend possible
Ampere (2020) cp.async copie globale → partagée sans passer par les registres
Hopper (2022) mbarrier attente asynchrone typée, avec compte d'octets
Hopper TMA un thread déclenche la copie d'un tenseur jusqu'à 5D
Hopper Clusters, DSMEM plusieurs blocs coopèrent sur des SM voisins
Hopper wgmma MMA asynchrone à l'échelle du warp group
Blackwell (2024) tcgen05 MMA sur paire de SM, opérandes en tensor memory
Blackwell Cluster Launch Control ordonnancement dynamique des clusters
CUDA 13.1 (2025) CUDA Tile modèle par tuiles, le compilateur place les threads

Ce que vous allez apprendre

  • Écrire un pipeline logiciel à plusieurs étages avec cp.async et mbarrier.
  • Utiliser TMA : descripteurs, mbarrier en mode expect-tx, multicast.
  • Faire coopérer plusieurs blocs via un cluster et la mémoire partagée distribuée.
  • Comprendre wgmma et tcgen05 assez pour lire du code CUTLASS.
  • Structurer un noyau en warps spécialisés (producteur / consommateur / épilogue), avec les schémas ping-pong et cooperative.
  • Situer CUDA Tile et savoir quand il remplace tout ce qui précède.

Ordre de lecture

Séquentiel. Chaque mécanisme s'appuie sur le précédent.

1 · Asynchronisme et mbarrier

cp.async, cuda::pipeline, mbarrier, et la construction d'un pipeline multi-étages. La brique de base de tout le reste.

2 · Le Tensor Memory Accelerator

Le mécanisme le plus important de Hopper. Descripteurs, copies multidimensionnelles, multicast vers un cluster, et pourquoi il change le calcul d'occupancy.

3 · Clusters et mémoire partagée distribuée

Le niveau de hiérarchie ajouté par Hopper : des blocs qui se synchronisent et lisent la mémoire partagée les uns des autres.

4 · wgmma et tcgen05

Les instructions matricielles asynchrones, leurs descripteurs, la tensor memory, et ce qu'il faut en savoir pour lire CUTLASS.

5 · La spécialisation des warps

Producteur / consommateur, les schémas ping-pong et cooperative, et le lien direct avec l'architecture des megakernels.

6 · CUDA Tile

Le modèle de programmation introduit en CUDA 13.1 : on décrit des tuiles, le compilateur place les threads. Ce qu'il apporte, ce qu'il ne fait pas encore.

Un avertissement

Ce chapitre décrit des mécanismes que peu de gens écrivent à la main

TMA, wgmma, tcgen05 et la spécialisation des warps sont extraordinairement pénibles à écrire directement. Les descripteurs sont des champs de bits, les dispositions de registres sont des tables, et une erreur donne un résultat faux sans message d'erreur.

En production, on utilise CUTLASS/CuTe, ThunderKittens, Triton ou Gluon, qui encapsulent tout cela.

Alors pourquoi ce chapitre ? Parce que ces bibliothèques exposent leurs abstractions dans ces termes, parce qu'on ne peut pas déboguer ce qu'on ne comprend pas, et parce que les megakernels sont l'un des rares endroits où l'on écrit encore ce niveau à la main.

Aucun code de cette partie n'a été exécuté

L'environnement de rédaction n'a ni GPU Hopper/Blackwell ni compilateur CUDA. Les extraits suivent les documentations PTX et CUDA citées. Les signatures exactes des intrinsèques évoluent d'une version de CUDA à l'autre : vérifiez toujours dans la documentation de votre toolkit.