Améliorer les performances des kernels CUDA avec le déversement des registres vers la mémoire partagée (CUDA 13.0)

Vue d’ensemble

Lorsqu’un noyau CUDA nécessite plus de registres matériels que ceux disponibles, le compilateur peut faire du spilling vers la mémoire locale, située dans la mémoire globale du dispositif. Ce spilling peut impacter les performances à cause de l’accès à la mémoire locale plus lent. CUDA 13.0 introduit une optimisation: le spilling des registres vers la mémoire partagée pour les noyaux CUDA. Le compilateur privilégie le spilling dans la mémoire partagée disponible et recourt à la mémoire locale si nécessaire. Cette mémoire on‑chip à faible latence rapproche les données déversées de la SM, réduisant la pression sur le cache L2 et améliorant les performances dans les chemins à forte pression de registres. Dans CUDA 13.0, PTXAS ajoute le support du spilling des registres vers la mémoire partagée pour les noyaux CUDA. Si cette fonctionnalité est activée, le compilateur tente d’abord de déverser dans la mémoire partagée et bascule vers la mémoire locale lorsque l’espace est insuffisant, ce qui garantit la correction du programme. Cette optimisation complète les caches L1/L2 plus grands et vise à réduire la latence des données stockées dans la mémoire on‑chip. Auparavant, tous les spills allaient vers la mémoire locale ; ce nouveau chemin peut générer des gains notables dans les kernels à forte pression de registres, comme les boucles ou les sections fréquemment exécutées. L’évaluation a été réalisée sur des kernels CUDA de la bibliothèque QUDA, utilisée pour les calculs Lattice QCD sur les GPUs, et a montré des gains typiques autour de 5–10%, attribués à la réduction ou l’élimination des spills vers la mémoire locale. L’optimisation est disponible uniquement avec CUDA 13.0 ou ultérieur. L’activation explicite via PTX inline est nécessaire.

Ce qui change dans le compilateur

Sous CUDA 13.0, le PTXAS introduit un mécanisme d’activation qui redirige les spills vers la mémoire partagée lorsque l’espace le permet. Lorsqu’elle est activée, la mémoire partagée est utilisée en priorité pour les spills et la mémoire locale n’est utilisée que si la mémoire partagée est épuisée. Un noyau compilé avec cette fonctionnalité peut afficher, dans les sorties Nsight Compute, l’utilisation de la mémoire partagée par bloc (par exemple, 46080 octets de smem). L’objectif est de maintenir les données déversées près de la SM et de réduire les évictions L2 dans les zones sensibles. Cette optimisation est activée en insérant une directive PTX (pragma) enable_smem_spilling via assembleur inline dans la fonction, juste après sa déclaration. La fonctionnalité est valable uniquement dans le cadre d’une fonction. Elle ne doit pas être utilisée lorsque les bornes de lancement ne sont pas explicitement spécifiées, car l’estimation du nombre de blocs par SM peut devenir inexacte et limiter la concurrence. Lorsque les bornes de lancement sont bien définies et l’utilisation de la mémoire partagée est stable, activer le spilling vers la mémoire partagée peut améliorer les performances.

Caractéristiques clés

• Support PTXAS pour le déversement des registres vers la mémoire partagée (CUDA 13.0+).
• Spill prioritaire vers la mémoire partagée; rebasculage vers la mémoire locale si l’espace est insuffisant.
• Améliorations potentielles des performances liées à la réduction de la latence et à une moindre pression sur la L2 dans les zones critiques.
• Activation via inline assembly PTX: enable_smem_spilling, placé après la déclaration de la fonction.
• Valide uniquement dans le cadre d’une fonction; recommandé lorsque les limites de lancement sont clairement définies.
• Illustrations de gains typiques dans des charges comme QUDA lattice QCD.
• Disponible uniquement dans CUDA 13.0 et versions ultérieures; les toolkits antérieurs ne supportent pas le spilling dans la mémoire partagée.
• Si la mémoire partagée par bloc est généralement sous-utilisée, l’optimisation peut quand même être bénéfique lorsque la pression de registres est élevée.

Cas d’utilisation courants

• Kernels avec une forte pression de registres et des boucles serrées où le spilling vers la mémoire locale était un goulot d’étranglement.
• Charges où une grande partie de la mémoire partagée par bloc reste inutilisée, rendant possible le stockage des spills dans la mémoire on‑chip.
• scénarios où l’occupation est limitée par les registres et non par la mémoire partagée.
• Applications avec des bornes de lancement bien définies et un usage prévisible de la mémoire partagée qui tolèrent l’augmentation du footprint par bloc.

Setup & installation (commandes exactes)

• Assurez‑vous d’utiliser CUDA 13.0 ou supérieur.
• Activez le spilling en mémoire partagée en insérant une pragma PTX en assembleur inline après la déclaration de la fonction: enable_smem_spilling. Cette activation est décrite comme valide dans le cadre d’une fonction et doit être utilisée lorsque les bornes de lancement sont explicitement définies. Note: la source décrit la méthode d’activation et les contraintes, mais ne fournit pas de commandes d’installation externes autres que l’exigence de version du toolkit CUDA.

// Activer le spilling en mémoire partagée (CUDA 13.0+)
// Placer après la déclaration de la fonction
// Inline PTX pragma: enable_smem_spilling

Quick start (exemple minimal exécutable)

Ci‑dessous deux exemples de noyaux minimaux pour illustrer l’idée conceptuellement. Le premier montre un noyau simple sans l’activation de spilling. Le second indique où et comment l’activation de spilling serait appliquée dans le code, selon les directives de CUDA 13.0.

// Noyau minimal (baseline, sans spilling)
extern "C" __global__ void demo_kernel(float* in, float* out) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float v = in[idx];
out[idx] = v * 2.0f;
}

// Noyau minimal (conceptuel: enable_smem_spilling après la déclaration de la fonction)
extern "C" __global__ void demo_kernel(float* in, float* out) {
// Activer spilling en mémoire partagée (CUDA 13.0+)
// Inline PTX pragma: enable_smem_spilling
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float v = in[idx];
out[idx] = v * 2.0f;
}

Remarque: la syntaxe exacte pour activer via inline PTX est décrite dans la source comme une pragma PTX placée après la déclaration de la fonction. Les exemples ci‑dessus illustrent l’emplacement et l’intention, et non une directive compilateur littérale.

Avantages et inconvénients

• Avantages
• Gains de performance potentiels (typique 5–10%) dans les kernels à forte pression de registres, en réduisant les spills vers la mémoire locale.
• Maintient les spills dans la mémoire partagée on‑chip, réduisant la latence et la pression L2 lorsque l’espace est disponible.
• Fournit une voie d’optimisation ciblée pour les kernels avec des limites de lancement bien définies et un usage stable de la mémoire partagée.
• Inconvénients
• Activation explicite requise; ce n’est pas automatique pour tous les kernels.
• Non applicable si la mémoire partagée par bloc est déjà saturée ou si la pression de registres n’est pas le goulot d’étranglement.
• Peut influencer l’occupation si les limites de lancement ne sont pas clairement définies, en raison d’estimations inexactes du nombre de blocs par SM.
• Disponible uniquement dans CUDA 13.0 et ultérieur; les toolkits plus anciens ne prennent pas en charge le spilling dans la mémoire partagée.

Alternatives (comparaisons rapides)

| Approche | Avantages | Inconvénients |---|---|---| | Spill en mémoire partagée (CUDA 13.0+) | Réduit les spills en mémoire locale; utilise la mémoire on‑chip | Activation explicite requise; dépend de la disponibilité de mémoire partagée |Spill vers mémoire locale (comportement pré‑CUDA 13.0) | Pas de modification de code nécessaire | Latence plus élevée; pression L2 accrue |Optimisation du code pour réduire la pression des registres | Réduit les spills sur l’ensemble des toolchains | Peut nécessiter une refonte importante; pas toujours faisable |

Pricing ou Licence

Aucune information sur les tarifs ou licences n’est fournie dans la source. La fonctionnalité est décrite comme partie de l’outil CUDA 13.0.

Références

• NVIDIA Dev Blog : How to Improve CUDA Kernel Performance with Shared Memory Register Spilling