Architecture GPU NVIDIA

Pourquoi ce chapitre conditionne tout le reste

On ne peut pas écrire un kernel CUDAGPUIAProcesseur graphique parallélisant massivement les calculs matriciels ; indispensable à l'entraînement et à l'inférence des modèles de deep learning.Voir dans le glossaire correct sans un modèle mental du matériel qui l'exécute. Les optimisations qui suivront dans cette formation — coalescing mémoire, mémoire partagée, streams, tuning d'occupancy — ne sont que des conséquences directes de l'architecture décrite ici. Un développeur qui code CUDA sans savoir ce qu'est un warp optimise à l'aveugle : il peut améliorer un kernel de 3 % en croyant avoir trouvé un levier majeur, ou dégrader silencieusement les performances en pensant faire mieux.

Ce chapitre ne survole pas la théorie GPU en général. Il se concentre sur ce qui change réellement le comportement d'un kernel : la façon dont un GPU NVIDIA découpe son silicium, ordonnance ses threads et expose (ou cache) ces détails dans la documentation constructeur.

Le Streaming Multiprocessor, unité de calcul réelle

Un GPU NVIDIA n'est pas un bloc de calcul monolithique. C'est une grille de Streaming Multiprocessors (SM) — de 20 à plus de 130 selon la puce — chacun étant une unité de calcul quasi autonome. Chaque SM embarque :

• un jeu de CUDA cores (unités arithmétiques scalaires, FP32/INT32 principalement) ;
• un ou plusieurs Tensor Cores (unités matricielles dédiées, depuis Volta) ;
• un ordonnanceur de warps (warp scheduler) qui sélectionne à chaque cycle quel warp émet une instruction ;
• un fichier de registres partagé entre tous les threads actifs sur le SM ;
• une mémoire partagée (shared memory) rapide, on-chip, configurable en partie comme cache L1.

Ce découpageChunkingIADécoupage d'un document en segments de taille fixe ou sémantique avant indexation vectorielle, pour optimiser la récupération RAG.Voir dans le glossaire a une conséquence directe sur la programmation : un bloc de threads est toujours exécuté en entier sur un seul SM, jamais réparti sur plusieurs. Le SM est donc l'échelle à laquelle raisonner pour tout ce qui touche à la contention de ressources — registres, mémoire partagée, occupancy.

CUDA cores et Tensor Cores : deux familles d'unités, deux usages

Les CUDA cores sont des unités arithmétiques scalaires classiques : elles traitent une opération FP32, FP64 ou INT32 par thread, par cycle. C'est le cheval de trait de tout kernel généraliste — tri, réduction, traitement d'image, simulation physique.

Les Tensor Cores, introduits avec Volta (2017) et largement étendus depuis, sont des unités matricielles : elles calculent directement des produits de matrices de petite taille (par exemple 4×4 ou 8×8 selon la génération et la précisionprécisionIAProportion des alertes émises par un modèle qui sont justifiées. Elle s'oppose au rappel : améliorer l'une dégrade l'autre.Voir dans le glossaire) en un nombre de cycles très inférieur à ce qu'il faudrait avec des CUDA cores enchaînant des multiplications-accumulations. Elles sont conçues pour l'algèbre linéaire dense qui domine l'entraînement et l'inférenceinférenceIAUtilisation d'un modèle déjà entraîné sur une donnée nouvelle. Peu coûteuse à l'unité mais répétée à chaque requête, elle constitue le coût récurrent d'exploitation.Voir dans le glossaire des réseaux de neuronesréseau de neuronesIAFonction mathématique composée de neurones artificiels organisés en couches, dont les coefficients sont ajustés pendant l'entraînement. Malgré son nom, il n'a presque rien de commun avec un cerveau.Voir dans le glossaire : convolutions et couches denses se ramènent presque toujours à des GEMM (General Matrix Multiply).

Les Tensor Cores ne s'activent pas automatiquement. Un kernel écrit en CUDA C++ classique, ou un appel BLAS mal configuré, peut tourner entièrement sur les CUDA cores sans jamais toucher les Tensor Cores — même sur un GPU qui en est équipé. Leur usage passe par des bibliothèques dédiées (cuBLAS, cuDNN, WMMA/MMA API) ou par des frameworks qui les ciblent explicitement (précision mixte FP16/BF16/TF32/FP8, formats de tenseurs alignés). Un profil qui montre un throughput FP32 proche du pic CUDA-core mais très en dessous du pic Tensor-Core annoncé sur la fiche produit est un signal que les Tensor Cores ne sont simplement pas sollicités.

Cette distinction est la première chose à vérifier quand un kernel d'IAintelligence artificielleIAEnsemble des techniques permettant à un programme d'accomplir une tâche qui demanderait de l'intelligence humaine. Le terme couvre aussi bien les systèmes à règles écrites que ceux qui apprennent de données.Voir dans le glossaire « n'atteint pas les performances annoncées » : la question n'est pas seulement combien de cœurs sont utilisés, mais lesquels.

Le warp : le vrai grain d'exécution

Le bloc et le thread sont des concepts du modèle de programmation. Le warp est un concept matériel, et c'est lui qui détermine la performance réelle. Sur toutes les architectures NVIDIA actuelles, un warp regroupe exactement 32 threads consécutifs d'un même bloc. Le warp scheduler d'un SM ne raisonne jamais thread par thread : il émet une instruction pour un warp entier, exécutée en lockstep sur les CUDA cores disponibles — c'est le modèle SIMT (Single Instruction, Multiple Threads).

Deux conséquences pratiques directes :

• Le nombre de threads par bloc devrait toujours être un multiple de 32. Un bloc de 100 threads occupe en réalité 4 warps (128 threads matériels), dont 28 threads du dernier warp sont inactifs mais consomment quand même un cycle d'émission.
• La divergence de branche coûte cher. Si un if (threadIdx.x < 16) sépare les threads d'un même warp en deux chemins différents, le matériel exécute les deux chemins l'un après l'autre, masquant à chaque passe les threads qui ne sont pas concernés (schéma ci-dessus). Un warp qui diverge systématiquement peut perdre jusqu'à 50 % de son débit d'instructions sur ce segment de code.

La divergence de branche ne se limite pas aux if/else explicites. Des boucles à nombre d'itérations variable selon le thread, des accès à des structures de donnéesdonnéesIAEnsemble d'informations structurées ou non utilisées pour entraîner, évaluer ou alimenter un modèle. La qualité, la quantité et la représentativité des données sont les facteurs décisifs pour les performances en apprentissage automatique.Voir dans le glossaire conditionnels, ou un simple return anticipé pour certains threads produisent le même effet. Le compilateur ne peut rien y faire : c'est une propriété de l'algorithme, pas du code généré.

Blocs et grilles : le mapping logiciel → matériel

Le modèle de programmation CUDA organise les threads en trois niveaux : thread, bloc, grille. Ce découpage logiciel se projette sur le matériel de façon précise :

Un SM peut héberger plusieurs blocs simultanément, tant que les ressources (registres, mémoire partagée, warps résidents) le permettent. C'est le GPU lui-même, via son ordonnanceur global, qui décide quel bloc va sur quel SM et quand — le développeur ne contrôle pas ce placement, seulement la taille des blocs et de la grille.

Le nombre maximal de threads par bloc est une limite matérielle (1024 sur la plupart des architectures récentes), pas un choix arbitraire. Dépasser cette limite dans la configuration de lancement d'un kernel produit une erreur au lancement, pas un ralentissement.

Occupancy : remplir le SM sans le saturer

L'occupancy est le ratio entre le nombre de warps effectivement résidents sur un SM et le nombre maximal de warps que ce SM peut héberger. Une occupancy de 100 % ne signifie pas une performance maximale : elle signifie seulement que le SM a de quoi masquer la latence mémoire en changeant de warp actif pendant qu'un autre attend une donnée.

Trois ressources limitent l'occupancy, et c'est presque toujours la plus contraignante des trois qui fixe le plafond réel :

• Registres par thread — un kernel gourmand en registres réduit le nombre de threads simultanément résidents, puisque le fichier de registres du SM est de taille fixe.
• Mémoire partagée par bloc — une allocation statique ou dynamique excessive limite le nombre de blocs corésidents.
• Nombre de threads par bloc — un bloc trop petit sous-utilise les warps schedulers ; un bloc trop grand limite le nombre de blocs simultanés.

L'outil nvcc --ptxas-options=-v affiche, pour chaque kernel compilé, le nombre de registres utilisés par thread et la mémoire partagée statique consommée. C'est le point de départ pour tout calcul d'occupancy théorique, avant même de lancer le profiler Nsight Compute. Une occupancy théorique correcte n'est pas une garantie de performance, mais une occupancy trop faible (moins de 25 % sur un kernel limité par la latence mémoire) est presque toujours un problème réel à corriger.

Un kernel de réduction utilise 64 registres par thread sur un SM qui en propose 65536 au total. Cela autorise au maximum 65536 / 64 = 1024 threads résidents, soit 32 warps — potentiellement l'occupancy maximale de ce SM. Le même kernel, une fois modifié pour dérouler une boucle et gagner en instructions par cycle, grimpe à 96 registres par thread : la résidence maximale tombe à 682 threads, soit 21 warps. Le gain d'instruction-level parallelism peut alors être annulé, voire dépassé, par la perte de capacité à masquer la latence mémoire. C'est un arbitrage à mesurer, jamais à supposer.

Générations récentes : ce qui change vraiment

Les fiches produit NVIDIA mettent en avant des TFLOPS crête qui ne se traduisent en gains réels que si le code cible explicitement les nouveautés de chaque génération.

Ce tableau ne remplace pas la fiche technique du modèle précis visé — les déclinaisons (SXM, PCIe, variantes mémoire) diffèrent nettement en bande passante et en consommation. Il donne l'axe d'évolution : à chaque génération, NVIDIA ajoute des formats de précision réduite pour les Tensor Cores et des mécanismes de communication inter-GPU, plus que des gains bruts de fréquence ou de nombre de cœurs.

Lire un spec sheet sans se faire piéger

Une fiche technique NVIDIA affiche généralement des TFLOPS « avec sparsité » ou en précision réduite (FP8, FP4) comme métrique phare. Ces chiffres sont crêtes théoriques, atteignables seulement dans des conditions précises : format de données aligné, sparsité structurée effectivement présente dans le modèle, kernel écrit pour exploiter ce mode.

Méthode pratique pour évaluer un GPU pour un projet donné :

1. Identifier la précision réellement utilisée dans le pipeline (FP32 d'entraînement classique, FP16/BF16 en mixed precision, INT8/FP8 en inférence quantifiée) et ne comparer que les TFLOPS de cette précision.
2. Vérifier si le chiffre affiché inclut la sparsité structurée — dans ce cas, diviser approximativement par deux pour obtenir l'équivalent dense, sauf si le modèle est effectivement élagué selon ce motif.
3. Croiser avec la bande passante mémoire (HBM) : un kernel memory-bound ne bénéficiera jamais du plein potentiel Tensor Core, quelle que soit la génération.
4. Regarder le nombre de SM et la taille du fichier de registres par SM, pas seulement le nombre total de « cœurs » — c'est ce qui borne l'occupancy réelle des kernels existants.
5. Vérifier NVLink/NVSwitch si le déploiement est multi-GPU : la topologie d'interconnexion pèse souvent plus que la puissance de calcul brute d'un seul GPU sur les charges d'entraînement distribué.

Pièges courants

• Comparer des TFLOPS de générations ou de précisions différentes comme s'ils étaient équivalents.
• Dimensionner des blocs à des tailles non multiples de 32, en pensant que le compilateur « arrondit » sans coût.
• Confondre nombre de CUDA cores et nombre de threads réellement exécutables en parallèle : le premier mesure le débit arithmétique crête, le second dépend de l'occupancy.
• Supposer qu'une occupancy à 100 % est toujours souhaitable, alors que certains kernels compute-bound performent mieux avec moins de warps mais plus de registres disponibles par thread (moins de spilling).
• Ignorer la divergence de branche parce que le code « compile et tourne » — il tourne, mais pas à la vitesse attendue.

Checklist avant d'optimiser un kernel

• La taille de bloc est un multiple de 32.
• Le nombre de registres par thread et la mémoire partagée par bloc ont été vérifiés (ptxas -v ou Nsight Compute).
• L'occupancy théorique a été calculée, pas supposée.
• Les branches conditionnelles dépendantes de threadIdx ont été identifiées et, si possible, réorganisées pour aligner la divergence sur les frontières de warp.
• La précision numérique cible (FP32/FP16/BF16/FP8) correspond à celle utilisée pour comparer les TFLOPS du spec sheet.
• Le kernel a été catégorisé compute-bound ou memory-bound avant de choisir le levier d'optimisation.

Cette base architecturale — SM, warps, occupancy, générations — est le socle sur lequel s'appuient les chapitres suivants consacrés au modèle mémoire, au coalescing et à l'écriture de kernels performants.