Hiérarchie mémoire et coalescing

Pourquoi la mémoire domine la performance CUDA

Un noyau 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 mal optimisé n'est presque jamais limité par le calcul brut : il est limité par la mémoire. Les architectures GPU modernes offrent plusieurs teraflops de puissance de calcul, mais cette puissance reste inutile si les unités de calcul attendent des 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 qui n'arrivent pas assez vite. Le ratio entre opérations arithmétiques et octets transférés — l'intensité arithmétique — détermine si un noyau est compute-bound ou memory-bound. Dans la grande majorité des cas rencontrés en production (traitement de tenseurs, réductions, opérations d'indexation), le noyau est memory-bound, et la performance se joue entièrement sur la façon dont les threads d'un warp accèdent à la mémoire.

Ce chapitre part du principe que vous savez déjà lancer un noyau et raisonner en blocs et warps. L'objectif ici est de comprendre où vivent les données, combien coûte chaque niveau de la hiérarchie, et quels motifs d'accès déclenchent des pénalités silencieuses — silencieuses parce que le code compile et s'exécute normalement, il tourne simplement deux à dix fois plus lentement qu'il ne le devrait.

Un noyau qui « fonctionne » n'est pas un noyau correct du point de vue performance. Sur GPU, un facteur 5x à 10x entre une implémentation naïve et une implémentation qui respecte la hiérarchie mémoire est la norme, pas l'exception. Ne validez jamais un noyau critique sans être passé par Nsight Compute au moins une fois.

Vue d'ensemble de la hiérarchie

Chaque niveau de mémoire sur un GPU représente un compromis entre latence, capacité et portée (par thread, par bloc, ou globale au device).

La règle générale : plus un niveau est proche du cœur de calcul, plus il est rapide et petit. Les registres sont gratuits en cycles mais rares ; la mémoire globale est abondante mais coûte quatre cents cycles ou plus par accès non mis en cache. Toute stratégie d'optimisation consiste à faire remonter les données réutilisées vers les niveaux rapides (registres, shared) et à minimiser les allers-retours vers la mémoire globale.

Le coalescing : faire tenir un warp dans une transaction

Quand les 32 threads d'un warp exécutent une instruction de lecture ou d'écriture en mémoire globale, le contrôleur mémoire ne traite pas 32 requêtes indépendantes. Il regroupe les adresses demandées et émet le moins de transactions de 128 octets possible (32 octets sur certaines architectures pour des segments plus fins, mais raisonnez en 128 octets pour un warp de 32 threads accédant à des mots de 4 octets). Un accès est dit coalescé lorsque les 32 threads accèdent à des adresses consécutives et alignées, ce qui permet de satisfaire tout le warp avec une seule transaction.

Un tableau float data[N] parcouru avec data[threadIdx.x + blockIdx.x * blockDim.x] produit un accès parfaitement coalescé : le thread 0 lit l'octet 0, le thread 1 lit l'octet 4, etc. Les 32 threads couvrent exactement 128 octets contigus, alignés sur la frontière de transaction. Une seule transaction suffit.

À l'inverse, un accès en data[threadIdx.x * stride] avec stride > 1 (typique d'un parcours par colonne d'une matrice stockée en row-major) éparpille les adresses. Chaque thread touche un segment différent, et le contrôleur peut devoir émettre jusqu'à 32 transactions de 128 octets pour ne récupérer que 128 octets utiles au total — un facteur 32 de gaspillage de bande passante.

Les cas les plus fréquents de désoptimisation :

• Parcours par colonne d'une matrice row-major. Transposez le pattern d'accès, ou stockez en column-major si le parcours dominant est par colonne.
• Structure of Arrays (SoA) vs Array of Structures (AoS). Un struct Particle { float x, y, z; } stocké en tableau d'AoS force chaque thread à sauter des champs non désirés. Passer en SoA (float x[N]; float y[N]; float z[N];) restaure le coalescing sur chaque champ.
• Décalage d'alignement. Un accès à data[idx + 1] sur toute une grille décale la fenêtre de lecture d'un élément, cassant l'alignement sur la frontière de 128 octets pour une partie des warps.

Pour vérifier le coalescing sans deviner, utilisez Nsight Compute et regardez la métrique l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum rapportée au nombre de requêtes. Le ratio idéal est proche de 4 secteurs de 32 octets par requête de warp complet (soit 128 octets). Un ratio nettement supérieur signale un accès dispersé.

Bank conflicts en mémoire partagée

La mémoire partagée est organisée en 32 banks entrelacés, chaque bank traitant un mot de 4 octets par cycle. Un accès sans conflit se produit quand chaque thread du warp lit ou écrit dans un bank différent, ou quand plusieurs threads lisent exactement la même adresse (mécanisme de broadcast, sans pénalité). Un conflit de bank survient quand plusieurs threads accèdent à des adresses différentes qui tombent dans le même bank : ces accès sont alors sérialisés, un par un.

Le cas classique est l'accès à un tableau 2D en mémoire partagée avec un stride égal à un multiple de 32. Par exemple, __shared__ float tile[32][32]; accédé par colonne (tile[threadIdx.x][k]) : les 32 threads du warp accèdent à des lignes différentes mais à la même colonne. Comme chaque ligne fait 32 floats = 128 octets = exactement la largeur des 32 banks, tous les threads tombent dans le même bank. C'est un conflit de bank à 32 voies, le pire cas possible : l'accès qui devrait prendre un cycle en prend 32.

Les conflits de bank ne provoquent aucune erreur, aucun warning à la compilation, et le noyau reste fonctionnellement correct. Ils se traduisent uniquement par une chute de performance mesurable au profilage. C'est l'un des pièges les plus coûteux en temps de debug parce que rien ne signale le problème sauf le chronomètre.

La solution standard est le padding : déclarer le tableau partagé avec une colonne supplémentaire, __shared__ float tile[32][33];. Ce décalage d'un mot casse l'alignement qui causait la collision systématique — chaque ligne commence désormais sur un bank différent, et l'accès par colonne se répartit naturellement sur les 32 banks. Le coût est de 4 octets gaspillés par ligne, négligeable face au gain.

Autres leviers pour limiter les bank conflicts :

• Préférer des accès par ligne (contigus) plutôt que par colonne quand la structure de l'algorithme le permet.
• Regrouper les threads qui accèdent à la même adresse pour bénéficier du broadcast plutôt que de le fragmenter.
• Sur les architectures récentes, envisager une largeur de bank configurable (4 ou 8 octets) via cudaDeviceSetSharedMemConfig si l'application manipule majoritairement des doubles.

Mémoire constante : le broadcast comme arme

La mémoire constante (64 Ko adressables, déclarée avec __constant__) dispose d'un cache dédié optimisé pour un motif précis : tous les threads d'un warp lisent la même adresse au même moment. Dans ce cas, la lecture est diffusée (broadcast) à l'ensemble du warp en un seul cycle après le premier chargement. C'est le mécanisme idéal pour des coefficients de convolution, des paramètres de kernel partagés, ou une table de lookup fixe.

À l'inverse, si les threads d'un warp lisent des adresses différentes en mémoire constante, les accès sont sérialisés — le cache constant n'est pas conçu pour la diversité d'adresses, contrairement au cache L1/L2 sur la mémoire globale. N'utilisez la mémoire constante que pour des données réellement identiques pour tout le warp, jamais comme substitut générique à la mémoire globale.

Pinned vs pageable : le goulot d'étranglement du transfert hôte-device

Avant même d'atteindre la mémoire globale du GPU, les données doivent transiter depuis la RAM du CPU via PCIe (ou NVLink sur les configurations qui le supportent). Le comportement diffère radicalement selon le type de mémoire hôte utilisé :

• Mémoire pageable (allouée par malloc ou new classique) : le système d'exploitation peut la déplacer ou la swapper à tout moment. Un cudaMemcpy depuis une zone pageable oblige le driver à effectuer une copie intermédiaire vers un tampon pinned interne avant le DMA réel, ajoutant une latence et limitant la bande passante.
• Mémoire pinned (page-locked), allouée avec cudaMallocHost ou cudaHostAlloc : verrouillée en RAM physique, elle permet un accès direct par le moteur DMA du GPU. La bande passante mesurée est typiquement deux à trois fois supérieure à celle obtenue avec de la mémoire pageable sur le même bus, et c'est un prérequis absolu pour les transferts asynchrones (cudaMemcpyAsync) qui se chevauchent avec le calcul.

N'allouez pas systématiquement tous vos buffers hôte en pinned : c'est une ressource rare (verrouillée, elle réduit la mémoire disponible pour le reste du système) et son allocation est plus lente qu'un malloc classique. Réservez le pinned aux buffers effectivement transférés en boucle vers le GPU — typiquement les buffers d'entrée/sortie d'un pipeline d'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, pas les structures de contrôle côté hôte.

Pour les pipelines qui alternent copie et calcul, combiner mémoire pinned et streams CUDA multiples permet de recouvrir le transfert PCIe par l'exécution d'un noyau sur un autre flux de données, masquant une bonne partie de la latence de transfert.

Unified Memory : ce qui casse en production

L'Unified Memory (cudaMallocManaged) simplifie le code en exposant un espace d'adressage unique entre CPU et GPU, avec migration automatique des pages. C'est un excellent outil de prototypage, mais son usage en production mérite des réserves précises.

L'Unified Memory ne supprime pas le coût du transfert, elle le déplace et le rend implicite. Une faute de page déclenchée par le GPU sur une page encore résidente côté CPU coûte largement plus cher qu'un cudaMemcpyAsync explicite et planifié, parce qu'elle interrompt l'exécution du warp fautif le temps de la migration.

Les points de vigilance concrets :

• Fautes de page (page faults) au premier accès. Sans prefetch explicite (cudaMemPrefetchAsync), chaque première touche d'une page par le GPU déclenche une migration à la demande, avec un coût de latence imprévisible qui apparaît comme un pic dans le profil d'exécution.
• Oversubscription mémoire. Sur les architectures qui l'autorisent, l'Unified Memory permet d'allouer plus de mémoire que la capacité physique du device, avec pagination automatique vers l'hôte. Utile pour dépasser temporairement la limite VRAM, mais catastrophique en performance si le working set réel dépasse durablement la VRAM disponible : le système entre alors dans un régime de thrashing.
• Contention multi-GPU. Sur un nœud à plusieurs GPU, une page migrée en boucle entre deux devices (accès concurrents alternés) génère un trafic de migration qui peut dominer le temps d'exécution total.
• Granularité de migration. Les pages migrées ont une granularité fixe (généralement 4 Ko à 2 Mo selon la génération) : un accès isolé à quelques octets peut déclencher la migration d'une page entière, gaspillant de la bande passante si le reste de la page n'est pas utilisé.

En production, l'usage recommandé de l'Unified Memory est hybride : l'utiliser pour la commodité de développement, puis remplacer les chemins chauds identifiés au profilage par une gestion mémoire explicite (cudaMalloc + cudaMemcpyAsync + streams), en conservant cudaMemPrefetchAsync et cudaMemAdvise pour les cas où l'Unified Memory reste pertinente (données accédées rarement, prototypage rapide, portabilité entre architectures hétérogènes).

Checklist d'audit mémoire avant mise en production

• Le noyau est-il memory-bound ou compute-bound ? Vérifiez le ratio arithmétique/octet avant d'optimiser au hasard.
• Les accès en mémoire globale sont-ils coalescés pour le pattern d'accès dominant (vérifié via Nsight Compute, pas par lecture de code) ?
• La structure de données est-elle en SoA plutôt qu'en AoS pour les champs réellement parcourus en parallèle ?
• La mémoire partagée utilise-t-elle un padding pour éviter les conflits de bank sur les accès par colonne ?
• Les transferts hôte-device répétés utilisent-ils de la mémoire pinned et des streams asynchrones ?
• L'Unified Memory est-elle utilisée uniquement là où sa flexibilité justifie le risque de fautes de page, avec prefetch explicite sur les chemins chauds ?
• La pression sur les registres par thread limite-t-elle l'occupancy de façon disproportionnée par rapport au bénéfice obtenu (vérifiable avec --ptxas-options=-v à la compilation) ?

Cette liste n'est pas exhaustive, mais elle couvre les causes les plus fréquentes d'un écart de performance entre un noyau qui « marche » et un noyau qui exploite réellement la bande passante mémoire disponible.