Glossaire¶
Environ 130 termes, par ordre alphabétique. Chaque entrée renvoie au chapitre qui la développe.
A¶
AoS (Array of Structures) — Disposition mémoire où chaque élément est une structure complète. Mauvaise pour le GPU : les accès à un champ sont espacés de la taille de la structure. Voir SoA. → Performance 2
Ampere — Architecture NVIDIA de 2020 (sm_80, sm_86). Introduit
cp.async, BF16, TF32, la sparsité structurée 2:4.
Atomique — Opération lisant, modifiant et écrivant une adresse sans interruption possible. Sur GPU, s'exécute dans le L2 pour la mémoire globale. → Modèle 7
AWQ (Activation-aware Weight Quantization) — Méthode de quantification protégeant les canaux identifiés comme importants d'après la magnitude des activations. → IA 5
B¶
Banc (bank) — Une des 32 subdivisions de la mémoire partagée, chacune de 4 octets. Deux threads visant le même banc à des adresses différentes créent un conflit. → Performance 2
Bande passante — Débit de transfert mémoire, en octets par seconde. 3,35 To/s sur H100, 8,0 To/s sur B200.
BF16 (bfloat16) — Format 16 bits avec 8 bits d'exposant (même étendue que FP32) et 7 de mantisse. Le format dominant en apprentissage. → Fondations 6
Blackwell — Architecture NVIDIA de 2024-2025 (sm_100 centre de données,
sm_120 grand public). Introduit tcgen05, la Tensor Memory, FP4, le Cluster
Launch Control.
Bloc (thread block, CTA) — Groupe de threads affecté à un seul SM,
partageant de la mémoire partagée et pouvant se synchroniser par
__syncthreads().
→ Fondations 3
Bulle mémoire — Période au démarrage d'un noyau où aucune requête mémoire n'est en vol, donc où la bande passante est inutilisée. → Megakernels 1
C¶
CCCL (CUDA Core Compute Libraries) — Thrust + CUB + libcu++. → Écosystème 7
CDNA — Architecture AMD pour le calcul (MI200, MI300, MI350). Wavefront de 64, Matrix Cores, MFMA.
Chiplet — Puce constituant une partie d'un GPU. Un B200 en a 2, un MI300X en a 8 (XCD). → Megakernels 6
Cluster — Niveau de hiérarchie introduit par Hopper : groupe de 1 à 8 blocs garantis co-résidents sur des SM du même GPC, pouvant se synchroniser et partager leur mémoire partagée. → CUDA moderne 3
Cluster Launch Control — Mécanisme Blackwell permettant à un cluster de
demander dynamiquement du travail au matériel plutôt que de le déduire de son
blockIdx.
Coalescence — Fusion des accès mémoire des 32 threads d'un warp en un minimum de transactions. Une coalescence parfaite utilise 100 % des octets transférés. → Performance 2
Compute capability — Numéro de version identifiant les capacités d'une architecture : 8.0 (A100), 9.0 (H100), 10.0 (B200).
Cooperative groups — API CUDA fournissant des abstractions de groupes de
threads : thread_block, thread_block_tile, cluster_group, grid_group.
→ Modèle 4
cp.async — Instruction Ampere copiant de la mémoire globale vers la
mémoire partagée sans passer par les registres.
→ CUDA moderne 1
CTA (Cooperative Thread Array) — Terme PTX pour un bloc.
CU (Compute Unit) — Équivalent AMD du SM.
CUB — Bibliothèque de primitives parallèles (réduction, scan, tri) à plusieurs niveaux : device, bloc, warp.
CuTe — Couche de CUTLASS formalisant les layouts comme des paires (forme, pas) composables algébriquement. → Écosystème 3
CuTe DSL — Interface Python de CUTLASS 4.x. FlashAttention-4 est écrit dedans.
CUDA Graph — Séquence de noyaux enregistrée puis rejouée, divisant le coût de lancement par 4 à 8. → Modèle 6
CUDA Tile — Modèle de programmation par tuiles introduit dans CUDA 13.1 : Tile IR, cuTile Python, et une implémentation C++ annoncée. → CUDA moderne 6
CUTLASS — Bibliothèque de gabarits C++ de NVIDIA pour construire des GEMM.
D¶
DeepEP — Bibliothèque de DeepSeek pour le dispatch all-to-all des MoE distribués. V2 utilise un back-end NCCL par défaut. → IA 6
Divergence — Situation où les threads d'un même warp prennent des chemins différents. Coûte la somme des branches. → Fondations 3
DSMEM (Distributed Shared Memory) — Mécanisme Hopper permettant à un bloc d'accéder à la mémoire partagée d'un autre bloc du même cluster.
E¶
Épilogue — Traitement appliqué au résultat d'une GEMM avant son écriture : biais, activation, quantification. Fusionné, il évite un aller-retour en HBM.
Événement (CUDA) — Marqueur inséré dans un flux, servant à mesurer le temps ou à exprimer une dépendance entre flux.
Événement (megakernel) — Objet de synchronisation entre tâches dans le
ttGraph de MPK.
→ Megakernels 4
F¶
FlashAttention — Algorithme calculant l'attention par blocs sans jamais matérialiser la matrice \(S \times S\), grâce au softmax en ligne. → IA 3
Fleet — Abstraction de tâches liées aux chiplets pour les megakernels sur GPU multi-puces (AMD, 2026).
Flux (stream) — File d'opérations CUDA exécutées dans l'ordre. Deux flux différents peuvent s'exécuter en parallèle.
FP4 / FP6 / FP8 — Formats flottants de 4, 6 et 8 bits. FP8 se décline en E4M3 (poids/activations) et E5M2 (gradients).
Fusion — Combinaison de plusieurs opérations en un seul noyau, divisant le trafic mémoire par le nombre d'opérations fusionnées. → IA 4
G¶
GEMM (General Matrix Multiply) — \(\mathbf{C} = \alpha\mathbf{A}\mathbf{B} + \beta\mathbf{C}\). L'opération centrale de l'apprentissage profond. → IA 2
GEMM groupée — Plusieurs GEMM de tailles différentes exécutées ensemble. Essentielle aux MoE. Supportée par cuBLAS depuis CUDA 13.1.
Gluon — Langage bas niveau de Triton, exposant layouts, mémoire partagée et spécialisation des warps. → Écosystème 2
GPC (Graphics Processing Cluster) — Groupe de SM partageant des ressources. Les clusters Hopper sont contraints à un GPC.
GQA (Grouped-Query Attention) — Variante où plusieurs têtes de requêtes partagent une tête de clés-valeurs. Réduit le cache KV sans réduire le calcul.
Grille (grid) — Ensemble des blocs d'un lancement de noyau.
Grid-stride loop — Boucle où chaque thread traite plusieurs éléments espacés du nombre total de threads. Préserve la coalescence et découple grille et problème. → Modèle 2
H¶
HBM (High Bandwidth Memory) — Mémoire empilée à côté du die, à très large bus. La mémoire principale des GPU de centre de données.
Helion — DSL Python de Meta, plus haut niveau que Triton, compilant vers Triton. → Écosystème 5
HIP — API AMD, quasi identique à CUDA, compilant pour AMD et NVIDIA.
Hopper — Architecture NVIDIA de 2022 (sm_90). Introduit TMA, les clusters,
DSMEM, wgmma, FP8.
I¶
ILP (Instruction-Level Parallelism) — Nombre d'instructions indépendantes en vol simultanément. Substitue au TLP pour cacher la latence. → Performance 3
Intensité arithmétique — \(I = W/Q\), opérations par octet transféré. Comparée à \(I_{\text{crit}} = P/B\), elle détermine le régime. → Performance 1
Interpréteur (sur GPU) — Boucle de dispatch exécutant une séquence d'instructions à l'intérieur d'un megakernel. → Megakernels 3
L¶
L1 / L2 — Caches. L1 par SM, non cohérent ; L2 partagé par tout le GPU, point de cohérence.
Latence — Temps entre l'émission d'une requête et l'arrivée du résultat. 400-800 cycles pour la mémoire globale.
LDS (Local Data Share) — Équivalent AMD de la mémoire partagée. 160 Ko par CU sur CDNA 4.
libcu++ — Implémentation CUDA de la bibliothèque standard C++ :
cuda::atomic, cuda::barrier, cuda::pipeline.
M¶
Marlin — Noyau W4A16 de référence, recouvrant la déquantification avec la GEMM. → IA 5
mbarrier — Objet de synchronisation Hopper en mémoire partagée, comptant
des octets et disposant d'une phase.
→ CUDA moderne 1
Megakernel — Noyau unique et persistant exécutant tout un programme tensoriel, avec des dépendances fines en mémoire globale. → Megakernels
Mémoire partagée (shared memory) — Mémoire rapide sur le SM, gérée explicitement, partagée par les threads d'un bloc. 227 Ko adressables sur Hopper/Blackwell.
Mémoire unifiée — Espace d'adressage commun CPU/GPU (cudaMallocManaged,
ou cohérence matérielle sur Grace Hopper, MI300A, Apple Silicon).
MFMA (Matrix Fused Multiply-Add) — Instruction matricielle AMD, synchrone, à l'échelle du wavefront.
MFU (Model FLOPs Utilization) — Fraction du pic matériel effectivement utilisée par les FLOP « utiles » du modèle.
Mirage MPK — Premier compilateur de megakernel (CMU et al., 2025). → Megakernels 4
MMA (Matrix Multiply-Accumulate) — Famille d'instructions de tensor core.
MoE (Mixture of Experts) — Architecture où seuls \(k\) experts sur \(E\) sont activés par jeton. → IA 6
Mojo — Langage de Modular visant à unifier la pile, du système au noyau.
MX (Microscaling) — Standards OCP de formats à échelle par blocs : MXFP8, MXFP6, MXFP4.
N¶
NCCL — Bibliothèque de collectives multi-GPU de NVIDIA.
Noyau (kernel) — Fonction __global__ exécutée sur le GPU.
Noyau persistant — Noyau dont les blocs restent résidents et consomment une file de travail. → Megakernels 2
NVFP4 — Variante NVIDIA de FP4 avec une échelle FP8 E4M3 par bloc de 16.
NVLink — Interconnexion GPU-GPU de NVIDIA. 900 Go/s sur H100.
NVSHMEM — Bibliothèque PGAS permettant la communication inter-GPU depuis l'intérieur d'un noyau. → IA 7
O¶
Occupancy — Rapport entre warps résidents sur un SM et le maximum matériel (64). Un moyen de cacher la latence, jamais un objectif. → Performance 3
P¶
Padding — Ajout d'éléments inutilisés pour éviter les conflits de banc.
[32][33] au lieu de [32][32].
Page (mémoire partagée) — Unité d'allocation dynamique de la mémoire partagée dans un megakernel. 16 Ko chez Hazy Research, 32 Ko chez MPK.
PagedAttention — Gestion du cache KV par pages non contiguës (vLLM).
PDL (Programmatic Dependent Launch) — Mécanisme Hopper permettant à un noyau de commencer son prologue avant la fin du précédent.
PTX — Assembleur virtuel NVIDIA, portable et documenté. Compilé en SASS par
ptxas.
Q¶
Quantification — Représentation des poids ou activations sur moins de bits. Notation W\(x\)A\(y\). → IA 5
Quantification de vagues (wave quantization) — Perte d'efficacité quand le nombre de blocs n'est pas un multiple des blocs par vague.
R¶
Réduction — Combinaison de \(n\) éléments en un seul par un opérateur associatif. \(O(\log n)\) étapes en parallèle. → Modèle 5
Registre — Mémoire la plus rapide, privée à un thread. 65 536 par SM, 255 maximum par thread.
Roofline — Modèle de performance : \(\text{perf} \le \min(P, I \times B)\). → Performance 1
Rubin — Architecture NVIDIA annoncée en production au CES 2026. HBM4, 288 Go par R100. Compatible Blackwell au niveau code.
S¶
SASS — Assembleur natif d'une architecture NVIDIA. Non documenté officiellement.
Scan (prefix sum) — Toutes les réductions préfixes. Primitive de l'allocation de taille variable.
Sentinelle — Valeur impossible (NaN) servant de signal de disponibilité, remplaçant les atomiques. Technique de Kog. → Megakernels 7
SIMT (Single Instruction, Multiple Thread) — Modèle d'exécution NVIDIA : 32 threads exécutent la même instruction, avec des masques d'activation.
SM (Streaming Multiprocessor) — Unité de calcul du GPU. 132 sur H100, 148 sur B200. → Fondations 2
SoA (Structure of Arrays) — Disposition où chaque champ est un tableau contigu. La bonne disposition sur GPU.
Softmax en ligne — Calcul du softmax par blocs avec correction rétroactive. Exact, pas approché. → IA 3
Spécialisation des warps — Attribution de rôles différents (loader, compute, storer) aux warps d'un bloc. → CUDA moderne 5
Spilling — Déversement de registres en « mémoire locale », qui est en réalité de la mémoire globale. Toujours coûteux dans une boucle chaude.
Sparsité 2:4 — Motif de parcimonie structurée exploité par les tensor cores depuis Ampere. Double le débit annoncé, rarement utilisé en pratique.
Stream-K — Stratégie d'ordonnancement de GEMM découpant selon \(K\) pour équilibrer la charge.
Swizzling — Permutation XOR des adresses en mémoire partagée éliminant les conflits de banc sans gaspiller d'espace.
SYCL — Standard Khronos de programmation hétérogène en C++ à source unique.
T¶
Tâche (task) — Unité de travail d'un SM dans un megakernel.
tcgen05 — Famille d'instructions matricielles de Blackwell, avec accumulateur en Tensor Memory et coopération sur paire de CTA. → CUDA moderne 4
Tensor core — Unité matérielle effectuant \(\mathbf{D} = \mathbf{A}\mathbf{B} + \mathbf{C}\) sur de petites matrices en une instruction. → Fondations 5
TF32 — Format 19 bits utilisés dans les tensor cores à partir d'opérandes FP32. Activé par défaut dans PyTorch, modifie silencieusement les résultats.
Thread — Unité d'exécution la plus fine. Sur GPU, une voie d'un warp.
ThunderKittens — Bibliothèque de tuiles 16×16 de Hazy Research, base des megakernels de Stanford.
TLP (Thread-Level Parallelism) — Parallélisme obtenu par le nombre de threads. Mesuré par l'occupancy.
TMA (Tensor Memory Accelerator) — Unité Hopper copiant des tenseurs jusqu'à 5D entre mémoire globale et partagée, sur ordre d'un seul thread. → CUDA moderne 2
TMEM (Tensor Memory) — Mémoire Blackwell de 256 Ko par SM, dédiée aux opérandes et accumulateurs des tensor cores. 16 To/s en lecture.
Triton — DSL Python de programmation par tuiles, back-ends NVIDIA, AMD, Intel. Généré par PyTorch Inductor.
ttGraph — Représentation de MPK : graphe de tâches au niveau des SM, où tâches et événements alternent.
Tuile (tile) — Sous-bloc d'un tenseur traité ensemble. L'unité de raisonnement de la programmation GPU moderne.
V¶
Vague (wave) — Ensemble des blocs résidents simultanément sur tout le GPU.
Volta — Architecture NVIDIA de 2017. Introduit les tensor cores et le compteur ordinal par thread.
W¶
Warp — Groupe de 32 threads exécutés ensemble. Unité d'ordonnancement et d'accès mémoire.
Warp group — 128 threads (4 warps), unité d'exécution de wgmma.
Wavefront — Équivalent AMD du warp. 64 threads sur CDNA, 32 sur RDNA.
wgmma — Instruction matricielle asynchrone de Hopper, à l'échelle du warp
group, avec opérandes en mémoire partagée. Atteint 95 % du pic contre 62,9 %
pour mma.
X¶
XCD (Accelerator Compute Die) — Chiplet d'un GPU AMD. Un MI300X en a 8, chacun avec son L2.
Notations mathématiques¶
| Symbole | Signification |
|---|---|
| \(B\) | bande passante mémoire (octets/s) |
| \(P\) | débit de calcul crête (opérations/s) |
| \(I\) | intensité arithmétique (opérations/octet) |
| \(I_{\text{crit}}\) | intensité critique, \(P/B\) |
| \(W\) | travail total (opérations) |
| \(Q\) | données transférées (octets) |
| \(\Theta\) | nombre de paramètres d'un modèle |
| \(b\) | taille de lot |
| \(S\) | longueur de séquence |
| \(d\) | dimension du modèle |
| \(h\) | nombre de têtes d'attention |
| \(d_h\) | dimension par tête, \(d/h\) |
| \(L\) | nombre de couches |
| \(\eta\) | efficacité de coalescence |
Retour : Annexes · Index du cours