Aller au contenu

4 · La hiérarchie mémoire

Le chapitre le plus rentable du document. Presque tous les problèmes de performance GPU sont des problèmes de mémoire, et presque toutes les solutions consistent à faire vivre une donnée un niveau plus haut.

4.1 Les six niveaux

Du plus rapide au plus lent, avec les ordres de grandeur pour un H100 :

Niveau Portée Taille Latence Bande passante
Registres 1 thread 256 Ko / SM ~1 cycle ~100 To/s (agrégé)
Mémoire partagée 1 bloc ≤ 227 Ko / SM ~20-30 cycles ~30 To/s (agrégé)
Cache L1 1 SM partagé avec la mém. partagée ~30 cycles —
Cache L2 tout le GPU 50 Mo ~200 cycles ~7 To/s
HBM (globale) tout le GPU 80 Go ~400-800 cycles 3,35 To/s
Mémoire hôte CPU + GPU To ~2-10 µs 64 Go/s (PCIe 5 ×16)

Sur Blackwell s'ajoute la Tensor Memory (256 Ko/SM, 16 To/s en lecture), réservée aux tensor cores et traitée au chapitre 5.

L'ordre de grandeur qui résume tout

Entre un registre et la HBM, il y a un facteur ~500 en latence. Entre la HBM et la mémoire hôte via PCIe, encore un facteur ~50 en bande passante. Chaque franchissement de niveau vers le bas est une décision coûteuse.

4.2 Les registres

Chaque thread a ses propres registres. Sur Hopper et Blackwell, un SM dispose de 65 536 registres de 32 bits (256 Ko), et un thread peut en utiliser au maximum 255.

L'arbitrage est direct :

\[ \text{warps résidents par SM} \le \left\lfloor \frac{65\,536}{32 \times \text{registres par thread}} \right\rfloor \]
Registres/thread Warps résidents max Occupancy (sur 64)
32 64 100 %
64 32 50 %
128 16 25 %
255 8 12,5 %

Le compilateur choisit l'allocation. On peut la contraindre :

__global__ void __launch_bounds__(256, 4)   // 256 threads/bloc, ≥4 blocs/SM
mon_noyau(...) { ... }

ou avec -maxrregcount=N à la compilation.

Le spilling

Si un noyau a besoin de plus de registres que la limite, le compilateur déverse (spill) les excédents en « mémoire locale » — qui n'a de local que le nom : c'est de la mémoire globale. Un spill dans une boucle interne peut diviser les performances par cinq.

On le détecte à la compilation :

nvcc -Xptxas -v mon_noyau.cu
# ptxas info : Used 168 registers, 24 bytes spill stores, 24 bytes spill loads

Toute ligne spill non nulle dans un noyau chaud mérite investigation. Le billet de Hazy Research sur le megakernel tensor-parallèle mentionne explicitement les register spills comme une des optimisations restantes.

4.3 La mémoire partagée

C'est la ressource la plus importante à maîtriser. Une mémoire gérée explicitement par le programmeur, physiquement sur le SM, partagée par tous les threads d'un bloc.

__global__ void exemple() {
    __shared__ float tuile[32][32];        // statique, taille connue à la compilation
    extern __shared__ float dynamique[];   // taille passée au lancement
    ...
}
// lancement avec allocation dynamique :
exemple<<<grille, bloc, taille_en_octets>>>();

Pourquoi elle existe

Parce que le GPU n'a pas assez de cache pour deviner votre schéma de réutilisation (chapitre 1, §1.4). La mémoire partagée est un cache que vous programmez vous-même.

Le motif canonique, qu'on retrouvera partout :

1. Tous les threads du bloc chargent coopérativement une tuile
   de la mémoire globale vers la mémoire partagée (accès coalescés)
2. __syncthreads()
3. Chaque thread lit plusieurs fois dans la tuile (accès rapides,
   schémas arbitraires)
4. __syncthreads()
5. Passer à la tuile suivante

C'est ce motif qui transforme une multiplication matricielle naïve (intensité arithmétique \(I = 1/4\)) en une version par blocs (\(I \approx T/4\) pour une tuile \(T \times T\)). Détail en IA · GEMM.

Les bancs

La mémoire partagée est découpée en 32 bancs de 4 octets, entrelacés :

adresse (en float) : 0   1   2   3  …  31  32  33  …
banc               : 0   1   2   3  …  31   0   1  …

Un warp accède à la mémoire partagée en un cycle si ses 32 threads touchent 32 bancs différents (ou tous la même adresse, auquel cas il y a diffusion). Sinon il y a conflit de banc et l'accès est sérialisé.

L'exemple canonique du conflit :

__shared__ float m[32][32];
float x = m[threadIdx.x][0];   // ← tous les threads lisent la colonne 0

m[i][0] est à l'adresse \(32i\), donc au banc \(32i \bmod 32 = 0\). Les 32 threads visent le même banc, à des adresses différentes : conflit à 32 voies, 32 cycles au lieu d'un.

Le remède standard, le padding :

__shared__ float m[32][33];    // ← une colonne de plus
float x = m[threadIdx.x][0];   // adresse 33i → banc (33i mod 32) = i mod 32

Chaque thread touche maintenant un banc distinct. Une colonne de 128 octets gaspillée achète un facteur 32.

Traité en détail dans Performance · Coalescence et conflits de banc.

Elle partage le budget avec le L1

Sur toutes les architectures depuis Volta, mémoire partagée et cache L1 sont le même silicium, partitionné dynamiquement. Sur H100/B200 : 228 Ko au total. Demander 227 Ko de mémoire partagée par bloc laisse ~1 Ko de L1 — parfois le bon choix, souvent non.

Au-delà de 48 Ko par bloc, il faut le demander explicitement :

cudaFuncSetAttribute(mon_noyau,
    cudaFuncAttributeMaxDynamicSharedMemorySize, 227 * 1024);

C'est un piège classique : sans cet appel, le lancement échoue silencieusement avec cudaErrorInvalidValue.

4.4 Les caches L1 et L2

Contrairement à la mémoire partagée, ils sont automatiques. On les influence mais on ne les programme pas.

L1 : par SM, non cohérent avec les autres L1. Sur Hopper/Blackwell il gère les lectures globales, les copies cp.async, et le cache de constantes.

L2 : partagé par tout le GPU, c'est le point de cohérence. Toute communication entre blocs passe par lui (ou par la HBM). 50 Mo sur H100, 126 Mo sur B200, en 2 partitions sur Hopper et 4 sur Blackwell.

Instructions utiles pour piloter le comportement :

Directive Effet
__ldg(ptr) lecture en lecture seule via le cache de textures
ld.global.nc idem, en PTX
.cs (cache streaming) évincer rapidement, pour les données lues une fois
.cg (cache global) ne pas mettre en L1, seulement en L2
cudaAccessPolicyWindow marquer une plage comme persistante en L2 (Ampere+)

Le hint « non temporel »

Pour un flux de poids lu une seule fois — le cas exact de l'inférence LLM — il est contre-productif de polluer le L1 et le L2. Les non-temporal hints disent au matériel « ne garde pas ça ». Le monokernel Kog sur MI300X les utilise explicitement pour son flux continu de poids.

4.5 La mémoire globale (HBM)

C'est là que vivent vos données. Trois caractéristiques déterminent tout.

Elle est loin. 400 à 800 cycles de latence. À 1,7 GHz, cela fait 250 à 470 ns. Pendant ce temps, un SM peut exécuter des milliers d'instructions — d'où le besoin de warps en vol.

Elle est large. Une transaction fait 32 octets minimum (parfois 64 ou 128 selon la granularité de cache). Lire 4 octets coûte 32 octets de trafic.

Elle est saturable. 3,35 To/s sur H100 est une limite dure. Aucune astuce logicielle ne la dépasse ; le seul moyen d'aller plus vite est de transférer moins.

La règle de coalescence

Si les 32 threads d'un warp lisent 32 float consécutifs et alignés sur 128 octets, le matériel émet 4 transactions de 32 octets et utilise 100 % des octets transférés.

int i = blockIdx.x * blockDim.x + threadIdx.x;
float x = donnees[i];              // ✓ parfaitement coalescé
float y = donnees[i * 2];          // ✗ 50 % d'efficacité
float z = donnees[permutation[i]]; // ✗ jusqu'à 12,5 %

Ce point est si important qu'il a son chapitre dédié.

Les accès vectorisés

Un thread peut lire 16 octets en une instruction :

float4 v = reinterpret_cast<const float4*>(donnees)[i];
// équivaut à 4 float, mais en 1 instruction LDG.128 au lieu de 4 LDG.32

Gain : moins d'instructions, moins de pression sur les unités load/store, et souvent 10 à 30 % sur un noyau limité par la mémoire. Contrainte : l'adresse de base doit être alignée sur 16 octets.

4.6 La mémoire hôte et les transferts

Le lien PCIe est le maillon faible : ~64 Go/s en PCIe 5.0 ×16, soit 52 fois moins que la HBM d'un H100.

Trois modes de mémoire hôte :

Mode Allocation Débit Usage
Paginable malloc ~6 Go/s à éviter
Épinglée (pinned) cudaMallocHost ~55 Go/s transferts explicites
Unifiée cudaMallocManaged variable prototypage, oversubscription

La mémoire paginable est lente parce que le pilote doit d'abord la copier dans un tampon épinglé interne. Toujours utiliser cudaMallocHost pour les transferts dans une boucle chaude.

La règle du transfert

Un transfert hôte→GPU d'un mégaoctet coûte ~18 µs en PCIe 5. Sur un H100, 18 µs suffisent à faire 1,2 × 10¹⁰ opérations flottantes. Si votre calcul est plus court que ça, il ne valait pas le transfert.

Corollaire : gardez les données sur le GPU aussi longtemps que possible. Enchaîner dix noyaux sur des données résidentes est bien meilleur qu'un aller-retour par étape.

Sur les plateformes à mémoire cohérente (Grace Hopper, Grace Blackwell, Apple Silicon, MI300A), cette contrainte s'atténue fortement : le CPU et le GPU partagent le même espace physique avec une cohérence matérielle.

4.7 Ce qu'un débutant doit faire de ce chapitre

Une méthode en trois questions, à appliquer à chaque noyau :

  1. Quelle donnée est lue plus d'une fois ? → la mettre en mémoire partagée.
  2. Les threads voisins lisent-ils des adresses voisines ? → sinon, réorganiser la disposition des données ou l'affectation des threads.
  3. Combien d'octets au total le noyau doit-il lire, au minimum ? → diviser par la bande passante de la carte : c'est votre plancher de temps d'exécution. Si vous en êtes à 3×, vous avez du travail ; si vous en êtes à 1,1×, arrêtez d'optimiser.

La troisième question est la plus sous-utilisée et la plus puissante. Elle transforme « mon noyau est-il rapide ? » en une comparaison chiffrée.

Résumé du chapitre

À retenir

  • Six niveaux : registres, mémoire partagée, L1, L2, HBM, hôte. Facteur ~500 de latence entre les extrêmes.
  • Les registres limitent l'occupancy ; les spills vont en mémoire globale et sont un désastre silencieux.
  • La mémoire partagée est un cache que vous programmez. 32 bancs de 4 octets ; les conflits sérialisent.
  • La mémoire globale exige des accès coalescés ; sinon on gaspille jusqu'à 7/8 de la bande passante.
  • PCIe est 50× plus lent que la HBM : gardez les données sur la carte.

Vérifiez que vous avez compris

Un noyau lit une matrice de 4 096 × 4 096 float et écrit sa transposée. Quel est le temps minimal sur un H100 ?

Volume lu : \(4096^2 \times 4 = 67{,}1\) Mo. Volume écrit : autant. Total 134,2 Mo.

\[t_{\min} = \frac{134{,}2 \times 10^6}{3{,}35 \times 10^{12}} = 40{,}1\ \mu s\]

Une transposition naïve (sans mémoire partagée) atteint typiquement 4 à 8× ce temps, parce que soit les lectures soit les écritures sont non coalescées. Une version par tuiles en mémoire partagée avec padding atteint 1,1 à 1,3×. C'est l'exercice canonique du parcours B.

Pourquoi __shared__ float tuile[32][32]; puis un accès tuile[threadIdx.y][threadIdx.x] ne pose-t-il pas de conflit de banc ?

Pour un warp donné, threadIdx.y est constant (les 32 premiers threads d'un bloc 32×32 ont tous y = 0) et threadIdx.x va de 0 à 31. L'adresse est donc \(32y + x\) avec \(x\) variant de 0 à 31 : les 32 threads touchent les 32 bancs distincts. Accès en un cycle.

C'est l'accès transposé, tuile[threadIdx.x][threadIdx.y], qui pose problème — et c'est justement celui dont a besoin une transposition, d'où le padding.

Vous constatez 48 bytes spill stores dans un noyau. Est-ce grave ?

Cela dépend d'où. 48 octets déversés dans du code de prologue exécuté une fois : négligeable. Les mêmes 48 octets dans une boucle interne exécutée 10 000 fois : 480 Ko de trafic mémoire globale supplémentaires par thread, ce qui peut dominer complètement le noyau.

Le profileur (Nsight Compute, section Memory Workload Analysis, ligne Local Memory) vous dit si le trafic local est significatif. Remèdes : réduire la taille des tableaux locaux, éviter l'indexation dynamique de tableaux en registres (elle force le spill), ou augmenter -maxrregcount.

Chapitre suivant : 5 · Les tensor cores

Sources de ce chapitre