Aller au contenu

2 · Anatomie d'un GPU

Du die entier jusqu'au registre individuel. Ce chapitre est descriptif : il donne les noms et les chiffres dont tout le reste du document se sert.

2.1 Vue d'ensemble : quatre niveaux

┌───────────────────────────────────────────────────────────┐
│  GPU (le die, ou les dies)                                │
│  ┌─────────────────────────────────────────────────────┐  │
│  │  GPC — Graphics Processing Cluster  (×8 sur B200)   │  │
│  │  ┌───────────────────────────────────────────────┐  │  │
│  │  │  SM — Streaming Multiprocessor                │  │  │
│  │  │  ┌─────────────┐ ┌─────────────┐              │  │  │
│  │  │  │ Partition 0 │ │ Partition 1 │  … (×4)      │  │  │
│  │  │  │ 32 cœurs    │ │ 32 cœurs    │              │  │  │
│  │  │  │ 1 tensor c. │ │ 1 tensor c. │              │  │  │
│  │  │  │ 1 ordonn.   │ │ 1 ordonn.   │              │  │  │
│  │  │  │ 64 Ko regs  │ │ 64 Ko regs  │              │  │  │
│  │  │  └─────────────┘ └─────────────┘              │  │  │
│  │  │  Mémoire partagée + L1 : 228 Ko (B200)        │  │  │
│  │  │  Tensor Memory : 256 Ko (B200 uniquement)     │  │  │
│  │  └───────────────────────────────────────────────┘  │  │
│  └─────────────────────────────────────────────────────┘  │
│  Cache L2 partagé  (50 Mo sur H100, 126 Mo sur B200)      │
│  Contrôleurs mémoire → HBM (80–192 Go)                    │
│  NVLink / PCIe → autres GPU, CPU                          │
└───────────────────────────────────────────────────────────┘

Retenez ce diagramme. La quasi-totalité des décisions d'optimisation consiste à choisir à quel niveau une donnée doit vivre.

2.2 Le SM, unité fondamentale

Le Streaming Multiprocessor est au GPU ce que le cœur est au CPU : l'unité qui possède des registres, exécute des instructions, et a sa propre mémoire rapide.

Un SM Hopper ou Blackwell contient :

  • 4 partitions (aussi appelées processing blocks), chacune avec :
    • 32 cœurs CUDA FP32/INT32,
    • 1 tensor core,
    • 1 ordonnanceur de warps (warp scheduler),
    • 1 unité de dispatch,
    • 16 384 registres de 32 bits (64 Ko) ;
  • une unité de fonctions spéciales (SFU) pour exp, log, sin, rsqrt — bien plus lente que les cœurs ordinaires, ce qui aura son importance dans FlashAttention-4 ;
  • des unités load/store (LSU) ;
  • 228 Ko de mémoire partagée + L1 combinés sur Blackwell (sm_100), 228 Ko également sur Hopper sm_90 avec 227 Ko adressables par un bloc ;
  • sur Blackwell uniquement, 256 Ko de Tensor Memory (TMEM), une mémoire distincte réservée aux opérandes et accumulateurs des tensor cores.

Pourquoi 4 partitions et pas 1 SM monolithique ?

Parce qu'un ordonnanceur qui doit choisir parmi 64 warps candidats à chaque cycle serait un circuit énorme et lent. En découpant en quatre, chaque ordonnanceur ne voit que 16 warps. C'est un compromis de complexité, et il a une conséquence pratique : un warp est affecté à une partition pour toute sa vie. Il ne peut pas migrer.

2.3 Le vocabulaire des « cœurs »

Il faut désamorcer une confusion marketing.

Quand une fiche technique annonce « 16 896 cœurs CUDA » pour un H100, cela signifie : 132 SM × 4 partitions × 32 unités FP32. Un « cœur CUDA » n'est pas un cœur au sens CPU : il n'a ni compteur ordinal indépendant, ni registres propres, ni possibilité d'exécuter une instruction différente de son voisin. Ce sont des voies (lanes) d'une unité SIMD de 32 voies.

La bonne analogie :

Terme GPU Équivalent CPU honnête
SM un cœur
Partition un port d'exécution SIMD
Cœur CUDA une voie d'un registre AVX-512
Warp (32 threads) une instruction SIMD 32 voies
Thread CUDA un élément d'un vecteur SIMD

Un H100 a donc, honnêtement, 132 cœurs très larges, pas 16 896 cœurs.

Erreur fréquente

Comparer « 16 896 cœurs GPU » à « 64 cœurs CPU » et en conclure un facteur 264. La bonne comparaison porte sur le débit d'opérations flottantes : un H100 fait ~67 TFLOPS en FP32 vectoriel, un CPU serveur haut de gamme ~5 TFLOPS. Le facteur réel est plutôt de 13 — et il monte à ~200 si l'on compare les tensor cores en BF16.

2.4 Chiffres réels, trois générations

A100 SXM H100 SXM B200 SXM
Architecture Ampere Hopper Blackwell
Cible de compilation sm_80 sm_90a sm_100a
SM 108 132 148
Mémoire partagée + L1 par SM 192 Ko 228 Ko 228 Ko
Adressable par bloc 163 Ko 227 Ko 227 Ko
Tensor Memory par SM — — 256 Ko
Registres par SM 256 Ko 256 Ko 256 Ko
Cache L2 40 Mo 50 Mo 126 Mo
Mémoire 80 Go HBM2e 80 Go HBM3 192 Go HBM3e
Bande passante 2,0 To/s 3,35 To/s 8,0 To/s
FP32 vectoriel 19,5 TFLOPS 67 TFLOPS ~80 TFLOPS
BF16 tensor (dense) 312 TFLOPS 990 TFLOPS ~2 250 TFLOPS
FP8 tensor (dense) — 1 979 TFLOPS ~4 500 TFLOPS
FP4 tensor (dense) — — ~9 000 TFLOPS
Threads résidents / SM 2 048 2 048 2 048
Dies 1 1 2 (liés par NV-HBI)

Les chiffres de SM pour A100/H100/B200 (108/132/148) sont ceux utilisés par le runtime de Mirage MPK, qui alloue 104/128/144 SM aux workers et 16 aux ordonnanceurs. Les valeurs de mémoire partagée Blackwell (228 Ko dont 227 Ko adressables) viennent du Blackwell Tuning Guide.

Attention aux chiffres « avec sparsité »

NVIDIA publie souvent des débits tensor core doublés grâce à la sparsité structurée 2:4 (deux zéros forcés sur quatre poids). Ce doublement suppose que votre modèle est effectivement élagué selon ce motif, ce qui est rare en pratique. Le tableau ci-dessus donne les valeurs denses. Un B200 annoncé à « 18 PFLOPS FP4 » fait 9 PFLOPS FP4 dense.

2.5 Blackwell est deux puces

Le B200 est physiquement deux dies reliés par une interconnexion appelée NV-HBI (10 To/s), présentés au programmeur comme un seul GPU CUDA. Cette fiction est excellente pour la portabilité et problématique pour la performance : deux SM sur des dies différents ne partagent pas leur L2.

Le B200 a d'ailleurs 4 partitions de L2, contre 2 sur Hopper (microbenchmarking Blackwell).

Le même problème existe, en plus prononcé, chez AMD : un MI300X est composé de 8 XCD (Accelerator Compute Dies) avec chacun son L2. C'est exactement ce que le papier Fleet et le monokernel de Kog exploitent : en dupliquant les données par die, on fait passer le taux de réussite du L2 de 12 % à 54 % sur certaines configurations.

À retenir

Depuis 2024, « un GPU » n'est plus une unité homogène. La localité de cache au niveau du die devient un paramètre d'optimisation, et le modèle de programmation CUDA/HIP ne l'expose pas encore directement.

2.6 Ce qui entoure les SM

Le cache L2

Partagé par tous les SM, il sert de tampon entre eux et la HBM. Sur H100 : 50 Mo, soit assez pour contenir les activations d'une couche de transformeur à petit lot, mais pas les poids. Il est non cohérent avec les caches L1 : si un SM écrit en L1 et qu'un autre lit en L2, il n'y a aucune garantie sans synchronisation explicite. C'est le fondement du problème que résolvent les megakernels.

Hopper introduit le contrôle de la persistance L2 (cudaAccessPolicyWindow) : on peut marquer une plage d'adresses comme « à garder en L2 », utile pour des poids réutilisés à chaque itération.

Les contrôleurs mémoire et la HBM

La mémoire à large bande passante (High Bandwidth Memory) est empilée verticalement à côté du die et connectée par un interposeur. Sa caractéristique : un bus très large (jusqu'à 8 192 bits) à fréquence modérée, plutôt qu'un bus étroit à haute fréquence comme la GDDR.

La conséquence pratique est fondamentale : la HBM délivre son débit uniquement sur des accès larges et contigus. Une transaction élémentaire fait 32 octets. Si vos 32 threads lisent 32 adresses éparpillées, vous déclenchez 32 transactions de 32 octets pour n'utiliser que 4 octets de chacune — vous obtenez 1/8 de la bande passante. C'est le problème de la coalescence.

Au-delà d'un GPU, les cartes sont reliées par NVLink (900 Go/s bidirectionnels par GPU sur H100, ~1,8 To/s sur B200) via des commutateurs NVSwitch. Un nœud DGX/HGX à 8 GPU est, du point de vue de la programmation, un seul espace mémoire adressable si l'on utilise NVSHMEM ou la mémoire symétrique. Voir IA · Multi-GPU.

2.7 Le côté AMD, en un tableau

Le vocabulaire diffère mais la structure est analogue.

Concept NVIDIA AMD
Unité de calcul SM CU (Compute Unit)
Groupe SIMD warp (32 threads) wavefront (64 threads sur CDNA)
Mémoire rapide sur puce shared memory LDS (Local Data Share)
Unité matricielle tensor core Matrix Core
Instruction matricielle mma, wgmma, tcgen05 MFMA (Matrix Fused Multiply-Add)
Langage CUDA HIP
Assembleur virtuel PTX (pas d'équivalent : GCN ISA direct)
Cible sm_90, sm_100 gfx942 (MI300X), gfx950 (MI350X)

Le MI350X/MI355X (CDNA 4, gfx950) a 256 CU, 160 Ko de LDS par CU, 288 Go de HBM3E à 8,0 To/s, et un support natif de MXFP8/MXFP6/MXFP4 avec échelle par blocs d'exposant.

La différence la plus lourde de conséquences est le wavefront de 64 au lieu du warp de 32 : tout code qui suppose 32 doit être audité lors d'un portage.

2.8 Combien de GPU différents faut-il connaître ?

Trois classes, et il faut savoir dans laquelle on se trouve.

Classe Exemples Ce qui les caractérise
Centre de données A100, H100, H200, B200, MI300X, MI355X HBM, FP64 rapide, NVLink, ECC, beaucoup de mémoire partagée
Station de travail / grand public RTX 4090, RTX 5090, L40S, RTX PRO 6000 GDDR (bande passante 3 à 8× plus faible), FP64 bridé, pas de NVLink sur le grand public
Embarqué / mobile Jetson Orin/Thor, GPU Apple, Adreno, Mali mémoire unifiée avec le CPU, contraintes d'énergie

Un noyau optimisé pour H100 est souvent mauvais sur RTX 4090 : moins de mémoire partagée par SM (100 Ko contre 227), pas de TMA, pas de clusters, bande passante 3,3× plus faible. C'est exactement le problème traité par Ada-MK, qui adapte les megakernels aux GPU Ada « à ressources contraintes ».

Résumé du chapitre

À retenir

  • Le SM est l'unité fondamentale : 4 partitions, chacune avec 32 voies FP32, un tensor core, un ordonnanceur et 64 Ko de registres.
  • Un « cœur CUDA » est une voie SIMD, pas un cœur.
  • Les trois chiffres à mémoriser pour votre carte : nombre de SM, mémoire partagée par SM, bande passante mémoire.
  • Depuis Blackwell et MI300X, un GPU est physiquement plusieurs puces ; la localité au niveau du die commence à compter.
  • Le vocabulaire AMD est différent mais la structure est la même, à l'exception majeure du wavefront de 64.

Vérifiez que vous avez compris

Combien de blocs de 256 threads peuvent résider simultanément sur un SM H100, si chaque bloc utilise 32 Ko de mémoire partagée ?

Deux contraintes :

  • threads : 2 048 / 256 = 8 blocs ;
  • mémoire partagée : 227 Ko / 32 Ko = 7 blocs (partie entière).

La contrainte la plus serrée gagne : 7 blocs. Il faudrait aussi vérifier les registres (2 048 threads × registres par thread ≤ 65 536 registres de 32 bits par SM) et la limite matérielle de blocs résidents par SM (32 sur Hopper).

Pourquoi lire 4 octets par thread dans un warp, à des adresses espacées de 128 octets, est-il catastrophique ?

Chaque thread touche une transaction de 32 octets différente. Le warp déclenche donc 32 transactions de 32 octets = 1 024 octets transférés, pour 128 octets utiles. Vous obtenez 12,5 % de la bande passante. Le même warp lisant 32 float contigus tient dans 4 transactions de 32 octets, soit 100 % d'efficacité.

Sur B200, la Tensor Memory fait 256 Ko par SM. À quoi sert-elle, alors qu'il y a déjà 228 Ko de mémoire partagée ?

À décharger le banc de registres. Avant Blackwell, les accumulateurs d'une MMA vivaient dans les registres des threads, ce qui limitait fortement la taille des tuiles (un accumulateur 128×128 en FP32 fait 64 Ko, soit un quart du banc de registres du SM). La TMEM est une mémoire dédiée, avec 16 To/s en lecture, ce qui permet à tcgen05.mma de traiter des tuiles 256×256 sans étrangler les registres. Détail au chapitre 5.

Chapitre suivant : 3 · Le modèle SIMT

Sources de ce chapitre