Aller au contenu

1 · Pourquoi un GPU existe

Le point de départ de tout : un GPU n'est pas « un CPU avec plus de cœurs ». C'est une machine qui résout un problème différent, et qui a accepté des compromis qu'un CPU refuse.

1.1 Le problème que résout un CPU

Un processeur central est conçu pour exécuter une chaîne d'instructions dépendantes le plus vite possible. C'est la contrainte du code ordinaire :

int x = a + b;
int y = x * c;      // dépend de x
if (y > seuil) { ... }   // dépend de y

Rien ici ne peut être fait en parallèle. La seule façon d'aller plus vite est de réduire le temps entre deux instructions dépendantes — c'est-à-dire de réduire la latence.

Un CPU moderne dépense la majorité de ses transistors à cela :

Mécanisme Ce qu'il achète Coût en transistors
Caches L1/L2/L3 de grande taille éviter d'aller en RAM (~80 ns) énorme
Prédiction de branchement ne pas vider le pipeline sur un if important
Exécution dans le désordre trouver du travail pendant qu'on attend énorme
Renommage de registres supprimer les fausses dépendances important
Préchargement matériel deviner la prochaine adresse modéré

Résultat : sur un die de CPU haut de gamme, moins de 10 % de la surface fait réellement de l'arithmétique. Le reste sert à ce que ces 10 % ne restent pas inactifs.

Intuition

Un CPU est un coursier à moto : il livre un colis très vite. Toute son ingénierie sert à ne jamais s'arrêter au feu rouge.

1.2 Le problème que résout un GPU

Le rendu d'une image, lui, ressemble à ceci :

pour chaque pixel de l'écran :          // 8 millions de pixels en 4K
    calculer sa couleur                 // indépendant des autres pixels

Il y a huit millions de tâches indépendantes. La latence de chacune n'a aucune importance : ce qui compte est le nombre de pixels finis par seconde, c'est-à-dire le débit.

Et quand on n'a pas besoin de réduire la latence, on peut supprimer tout le matériel qui servait à ça — et le remplacer par de l'arithmétique.

À retenir

CPU : minimiser le temps d'une tâche. GPU : maximiser le nombre de tâches par seconde. Toutes les autres différences en découlent.

1.3 Cacher la latence au lieu de la réduire

C'est le mécanisme central, et celui que les débutants comprennent en dernier.

Un GPU n'a pas d'exécution dans le désordre. Quand un thread demande une valeur en mémoire globale et doit attendre ~500 cycles, il ne se passe rien d'intelligent pour lui. Mais l'ordonnanceur bascule instantanément sur un autre groupe de threads prêt à s'exécuter.

Représentons quatre groupes de threads sur une même unité d'exécution (M = attente mémoire, C = calcul) :

temps →

groupe 1 :  CCC MMMMMMMMMMMM CCC MMMMMMMMMMMM
groupe 2 :      CCC MMMMMMMMMMMM CCC MMMMMMMM
groupe 3 :          CCC MMMMMMMMMMMM CCC MMMM
groupe 4 :              CCC MMMMMMMMMMMM CCC

unité    :  111 222 333 444 111 222 333 444 …   ← jamais inactive

Aucun groupe n'attend moins longtemps qu'avant. Mais l'unité de calcul, elle, travaille en permanence. La latence n'a pas disparu : elle a été recouverte par du parallélisme.

Sur un H100, un SM peut avoir 2 048 threads résidents — soit 64 warps de 32 threads. Le changement de contexte entre warps est gratuit : les registres de tous les warps résidents coexistent physiquement dans le banc de registres. C'est la raison pour laquelle ce banc est si grand (256 Ko par SM) et pourquoi consommer trop de registres par thread réduit le nombre de warps résidents.

Erreur fréquente

« Mon noyau est lent, je vais réduire le nombre d'accès mémoire par thread. » Peut-être. Mais si vous avez assez de warps en vol, ces accès sont déjà gratuits. La vraie question est : est-ce que la bande passante est saturée ? Voir le modèle roofline.

1.4 Le raisonnement en transistors

Voici l'arbitrage explicite, tel qu'il est fait par les concepteurs de puces.

Un budget de transistors est fixé par la surface de silicium et le procédé de gravure. On peut le dépenser en :

  • unités arithmétiques (ALU, FPU, tensor cores) → augmente \(P\), le débit crête ;
  • caches et logique de contrôle → réduit la latence effective.

Le CPU choisit la seconde option parce que ses charges de travail n'offrent pas assez de parallélisme pour cacher quoi que ce soit. Le GPU choisit la première parce que ses charges de travail en débordent.

Une conséquence directe, et souvent contre-intuitive :

\[ \text{cache par thread}_{\text{GPU}} \ll \text{cache par thread}_{\text{CPU}} \]

Sur un H100 : 50 Mo de L2 pour un maximum de 132 × 2 048 = 270 336 threads résidents, soit environ 190 octets de L2 par thread. Sur un CPU serveur, on est à plusieurs mégaoctets par thread. Le GPU n'a structurellement pas la place de mettre vos données en cache. Il faut donc que vous organisiez la réutilisation, explicitement, en mémoire partagée.

C'est là toute la difficulté du métier.

1.5 Une brève histoire, parce qu'elle explique le vocabulaire

Période Étape Ce qu'il en reste aujourd'hui
1999-2001 Pipeline graphique fixe (GeForce 256, « GPU » comme terme marketing) le mot « GPU »
2001-2006 Shaders programmables (vertex, pixel) — on détourne le rendu pour faire du calcul (« GPGPU ») les textures comme mécanisme de lecture, encore visibles dans __ldg
2006 G80 / CUDA 1.0 : architecture unifiée, langage C, mémoire partagée tout le modèle de programmation actuel
2008-2016 Fermi, Kepler, Maxwell, Pascal : FP64, caches, NVLink ECC, la hiérarchie de cache
2017 Volta : tensor cores, compteur ordinal par thread les tensor cores, la synchronisation explicite (__syncwarp)
2020 Ampere : cp.async, sparsité structurée, TF32 copies asynchrones
2022 Hopper : TMA, clusters, wgmma, FP8 la programmation asynchrone par tuiles
2024-2025 Blackwell : tcgen05, tensor memory, FP4 le modèle actuel
2026 Rubin : HBM4, en production depuis le CES 2026 compatible Blackwell au niveau code

Deux choses à retenir de cette liste.

D'abord, le modèle de programmation de 2006 est encore celui de 2026. Un noyau CUDA écrit pour un G80 compile et tourne sur un B200. C'est un exploit d'ingénierie de compatibilité, et c'est aussi la raison pour laquelle CUDA est difficile à déloger.

Ensuite, les nouveautés récentes sont toutes des mécanismes d'asynchronisme. De Ampere à Blackwell, l'histoire est : « faire en sorte que le déplacement de données ne bloque plus le calcul ». C'est le sujet de la partie 4.

Sur Rubin

NVIDIA a annoncé au CES 2026 que Rubin est en production, avec 288 Go de HBM4 par R100 et une disponibilité partenaire au second semestre 2026. Le point important pour ce document : le code Blackwell tourne sur Rubin. Les chiffres de ce cours restent valables ; les instructions spécifiques à Rubin apparaissent progressivement dans les PTX ISA livrés avec CUDA 13.4 (notes de version).

1.6 Ce qu'un GPU fait mal

Puisque tout découle du compromis débit/latence, on peut prédire les échecs.

Situation Pourquoi c'est mauvais sur GPU
Peu de travail (quelques milliers d'éléments) le lancement du noyau coûte plus que le calcul
Dépendances séquentielles longues rien à recouvrir, la latence apparaît nue
Branchements très divergents entre voisins les threads d'un warp s'attendent (chapitre 3)
Accès mémoire aléatoires et fins on paie une transaction de 32 octets pour lire 4
Récursion profonde, structures chaînées pile limitée, pointeurs = accès aléatoires
Latence de bout en bout critique (< 10 µs) le simple transfert PCIe coûte plus

Ces cas sont détaillés, avec les contournements possibles, dans Domaines · Ce qui ne va pas sur GPU.

Résumé du chapitre

À retenir

  • Un CPU minimise la latence, un GPU maximise le débit. Le reste suit.
  • Le GPU cache la latence par le parallélisme au lieu de la réduire ; le changement de contexte entre warps est gratuit parce que tous leurs registres coexistent.
  • Il y a structurellement très peu de cache par thread (~190 octets de L2 sur H100) : la réutilisation des données est votre travail.
  • Le modèle de programmation n'a pas changé depuis 2006 ; ce qui change, c'est l'asynchronisme du déplacement de données.

Vérifiez que vous avez compris

Un noyau fait 1 000 additions par thread sur des données déjà en registres, sans accès mémoire. Doubler le nombre de warps résidents l'accélère-t-il ?

Non, ou marginalement. Il n'y a pas de latence mémoire à cacher : les unités de calcul sont déjà saturées. Ajouter des warps ne fait que se disputer les mêmes unités. C'est un noyau limité par le calcul, et le seul levier est de réduire le nombre d'opérations ou d'utiliser des unités plus rapides (tensor cores, précision réduite).

Pourquoi un GPU n'a-t-il pas d'exécution dans le désordre, alors que c'est très efficace sur CPU ?

Parce qu'elle coûte extrêmement cher en transistors et en énergie, et qu'elle sert à trouver du travail indépendant dans un seul flux d'instructions. Le GPU a déjà des milliers de flux indépendants : il lui suffit de commuter entre eux, ce qui est infiniment moins cher. C'est le même service rendu par un mécanisme mille fois plus simple.

Sur H100, 132 SM × 2 048 threads = 270 336 threads. Ma matrice fait 1 024 × 1 024 éléments. Est-ce assez de parallélisme ?

1 024 × 1 024 = 1 048 576 éléments, soit ~3,9 threads résidents disponibles par élément si un thread traite un élément. C'est confortable en nombre de threads. Mais la vraie question est différente : combien d'octets transférez-vous par opération ? Une addition élémentaire sur des float lit 8 octets et écrit 4 pour 1 opération, soit \(I = 1/12\) — massivement limité par la mémoire. Le parallélisme est suffisant ; c'est la bande passante qui décidera.

Chapitre suivant : 2 · Anatomie d'un GPU

Sources de ce chapitre