Aller au contenu

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