Profiling avec Nsight

Pourquoi profiler avant d'optimiser

Optimiser 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 sans 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 de profiling revient à régler un moteur les yeux bandés. L'intuition sur ce qui limite les performances — trop de registres, accès mémoire mal alignés, occupation insuffisante — se trompe plus souvent qu'elle n'a raison, y compris chez des développeurs expérimentés. La chaîne d'outilstool callingIACapacité d'un agent ou d'un LLM à invoquer des outils externes (API, calcul, recherche) pendant le raisonnement.Voir dans le glossaire NVIDIA distingue deux niveaux d'analyse qui répondent à deux questions différentes. Nsight Systems répond à « où mon application perd-elle du temps, à l'échelle du programme entier ? ». Nsight Compute répond à « pourquoi ce kernel précis est-il lent ? ». Confondre les deux, ou sauter directement à l'optimisation fine sans vue d'ensemble, est la première cause de sessions de profiling improductives : on passe des heures à réduire l'occupancy d'un kernel qui ne représente que 2 % du temps d'exécution total.

Toujours commencer par Nsight Systems, jamais par Nsight Compute. La timeline applicative révèle souvent que le vrai problème n'est pas un kernel lent mais une synchronisation superflue, un transfert H2D redondant, ou une section CPU qui bloque le GPU en attente de données. Optimiser un kernel qui n'est pas sur le chemin critique ne change rien au temps d'exécution perçu.

Nsight Systems : la vue d'ensemble en timeline

Nsight Systems (nsys) capture l'exécution complète d'une application : threads CPU, appels à l'API CUDA, exécution des kernels par stream, transferts mémoire, et — si le code est instrumenté — plages NVTX personnalisées. La commande de base :

nsys profile --trace=cuda,nvtx,osrt -o rapport ./mon_application
nsys stats rapport.nsys-rep

Le fichier .nsys-rep s'ouvre dans l'interface graphique Nsight Systems, qui affiche une timeline multi-lignes : une ligne par thread CPU, une ligne pour la file des appels API CUDA, une ligne par stream GPU avec les kernels et copies mémoire qui s'y exécutent. C'est cette vue qui permet de répondre aux questions structurantes : le GPU est-il occupé en continu, ou y a-t-il des trous ? Les streams se recouvrent-ils réellement, ou s'exécutent-ils en série faute de dépendances mal gérées ? Le temps CPU entre deux lancements de kernel dépasse-t-il le temps d'exécution du kernel lui-même ?

Un trou dans la timeline GPU signale presque toujours l'une de ces causes : une synchronisation bloquante (cudaDeviceSynchronize, cudaStreamSynchronize) qui force le GPU à attendre le CPU, une dépendance de données non exprimée en streams multiples, ou un lancement de kernel dont la préparation côté CPU (calcul d'indices, allocation, cudaMemcpy synchrone) devient le facteur limitant.

Instrumentez le code source avec des plages NVTX (nvtxRangePushA("assemblage_matrice") / nvtxRangePop()) autour des sections logiques. Ces plages apparaissent comme des bandes colorées dans la timeline et permettent de faire correspondre directement une portion de code à une durée mesurée, sans deviner à partir des seuls noms de kernels CUDA.

Nsight Systems fournit aussi une vue statistique (nsys stats) qui agrège le temps total par type de kernel, par type d'appel API, et par type de transfert mémoire. C'est le point de départ pour prioriser : un kernel qui cumule 60 % du temps GPU mérite l'attentionattentionIAMécanisme par lequel un modèle pondère l'importance de chaque token du contexte lorsqu'il en traite un autre, quelle que soit la distance qui les sépare.Voir dans le glossaire de Nsight Compute ; un kernel à 1 % ne la mérite pas, même s'il « semble » mal écrit à la lecture du code.

Nsight Compute : diagnostic kernel par kernel

Une fois le ou les kernels critiques identifiés, ncu (Nsight Compute) permet une analyse fine d'une exécution isolée :

ncu --set full --kernel-name mon_kernel --launch-count 1 -o rapport_kernel ./mon_application

Le rapport se structure en sections. La section GPU Speed of Light donne en un coup d'œil deux pourcentages : l'utilisation du débit de calcul (compute throughput) et l'utilisation de la bande passante mémoire (memory throughput), chacun exprimé en proportion du maximum théorique de la carte. C'est la première chose à lire : si les deux pourcentages sont bas simultanément, le kernel n'est ni borné mémoire ni borné calcul — il est probablement limité par la latence (dépendances entre instructions, occupancy insuffisante) plutôt que par le débit.

La section Occupancy compare l'occupancy théorique (fonction du nombre de registres par thread, de la mémoire partagée utilisée, et de la configuration de lancement) à l'occupancy atteinte réellement pendant l'exécution. Un écart important entre les deux indique un déséquilibre de charge entre blocs, ou une exécution trop courte pour saturer les unités de calcul.

La section Warp State Statistics détaille les raisons pour lesquelles les warps sont en attente plutôt qu'actifs : Long Scoreboard (attente d'une donnée en mémoire globale), Barrier (attente à un __syncthreads()), Short Scoreboard (attente d'une opération mémoire partagée ou registre), MIO Throttle (unité mémoire saturée). C'est souvent l'information la plus actionnable du rapport : elle indique littéralement ce que les warps attendent le plus.

Une occupancy théorique maximale (100 %) n'est pas une garantie de performance. Un kernel à faible occupancy mais avec un excellent recouvrement calcul/mémoire grâce à l'instruction-level parallelism peut surpasser un kernel à occupancy maximale mais mal équilibré. Ne réglez jamais uniquement sur ce chiffre : croisez-le systématiquement avec le débit mémoire et la répartition des stalls de warp.

Le modèle roofline : borné mémoire ou borné calcul

Le modèle roofline relie la performance atteignable (en FLOP/s) à l'intensité arithmétique d'un kernel, définie comme le nombre d'opérations flottantes effectuées par octet transféré depuis la mémoire globale :

intensité_arithmétique = FLOPs_total / octets_transférés_total

Sur un graphe où l'axe horizontal est l'intensité arithmétique (échelle logarithmique) et l'axe vertical la performance atteinte, deux plafonds définissent la région de performance possible : un plafond horizontal fixé par le débit de calcul crête de la carte (indépendant de l'intensité), et un plafond diagonal fixé par la bande passante mémoire crête (proportionnel à l'intensité, puisqu'à bande passante fixe, plus on fait de calcul par octet, plus on peut atteindre de FLOP/s). Un kernel se place sous l'un ou l'autre plafond selon son profil : s'il est proche du plafond diagonal, il est borné mémoire — augmenter la bande passante effective (coalescing, réutilisation via mémoire partagée, réduction du volume de données) l'améliore ; augmenter le débit de calcul ne sert à rien. S'il est proche du plafond horizontal, il est borné calcul — l'inverse s'applique.

Un kernel de réduction (somme d'un grand tableau) lit un octet pour effectuer une seule addition : son intensité arithmétique est proche de 0,25 FLOP/octet en simple 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. Il colle au plafond mémoire quelle que soit la qualité du code : la seule optimisation utile est de maximiser la bande passante effective (coalescing parfait, utilisation de la mémoire partagée pour réduire les accès globaux redondants). À l'inverse, un GEMM dense bien écrit atteint des dizaines de FLOP par octet grâce au tuilage en mémoire partagée : il colle au plafond calcul, et seule une meilleure utilisation des unités de calcul (Tensor Cores, réduction de la divergence) l'améliore.

Nsight Compute calcule directement l'intensité arithmétique et positionne le kernel sur un graphe roofline dans la section Memory Workload Analysis combinée à Compute Workload Analysis — inutile de le calculer à la main pour un diagnostic rapide, mais comprendre le principe évite de mal interpréter un pourcentage de débit mémoire élevé comme un problème alors qu'il s'agit simplement de la limite physique du kernel.

Débogage avec cuda-gdb

Lorsqu'un kernel produit un résultat incorrect ou plante, cuda-gdb étend GDB au contextefenêtre de contexteIAQuantité de texte qu'un modèle peut prendre en compte simultanément : question, documents fournis et historique. Au-delà, les éléments les plus anciens sortent du champ.Voir dans le glossaire GPU : points d'arrêt dans le code device, inspection des variables par thread, navigation entre warps et blocs.

cuda-gdb ./mon_application
(cuda-gdb) break mon_kernel
(cuda-gdb) run
(cuda-gdb) cuda thread
(cuda-gdb) cuda thread (12,0,0)
(cuda-gdb) print mon_tableau[threadIdx.x]
(cuda-gdb) cuda kernel block thread

La commande cuda thread affiche le thread courant ; on peut basculer explicitement vers un thread ou un bloc donné pour inspecter son état au moment du breakpoint. C'est particulièrement utile pour un bug qui ne se manifeste que sur certains indices (effet de bord d'un accès hors limites, division par un cas particulier).

cuda-gdb nécessite une compilation avec -G -g (symboles de debug device, optimisations désactivées). Le comportement observé sous cuda-gdb peut différer du binaire optimisé de production : un bug de race condition peut disparaître (timing modifié) ou au contraire apparaître (l'optimiseur avait réordonné des accès d'une façon qui masquait le problème). Ne concluez jamais qu'un bug est résolu uniquement parce qu'il ne se reproduit plus sous debug — validez avec compute-sanitizer sur le binaire de release.

compute-sanitizer : détecter les erreurs mémoire et de synchronisation

compute-sanitizer (qui remplace l'ancien cuda-memcheck) regroupe plusieurs outils d'analyse dynamique, chacun ciblant une classe de bug :

memcheck est le plus utilisé au quotidien : un accès hors limites en mémoire globale ne provoque pas toujours un plantage immédiat — il peut corrompre silencieusement une zone mémoire voisine et ne se manifester que bien plus tard, par un résultat faux à un endroit sans lien apparent. racecheck est indispensable dès qu'un kernel utilise de la mémoire partagée sans synchronisation explicite entre écriture et lecture par des threads différents : ce type de bug est souvent non déterministe et invisible en test occasionnel, mais racecheck le détecte de façon fiable en instrumentant chaque accès.

Le surcoût de compute-sanitizer est important — un facteur 10 à 30 selon l'outil et le kernel, racecheck étant généralement le plus coûteux. Ne l'exécutez jamais sur un workload de production ou sur un jeu de données de taille réelle : réduisez la taille du problème pour la durée de l'analyse, et intégrez-le en CI sur des cas de test représentatifs plutôt qu'en profiling de performance.

Workflow de profiling recommandé

Une séquence qui évite le gaspillage de temps :

1. Nsight Systems sur l'application complète, jeu de données réaliste, pour identifier le ou les kernels/sections qui dominent le temps total.
2. Vérifier l'absence de trous anormaux dans la timeline GPU (synchronisations superflues, transferts non recouverts).
3. Nsight Compute ciblé sur les kernels identifiés en (1), en isolant un seul lancement (--launch-count 1) pour un rapport lisible.
4. Lire dans l'ordre : Speed of Light, puis Memory/Compute Workload Analysis pour situer le kernel sur le roofline, puis Warp State Statistics pour la cause de stall.
5. Appliquer une optimisation ciblée en fonction du diagnostic, puis rejouer nsys et ncu pour mesurer l'effet réel — jamais se fier à l'intuition sur le gain obtenu.
6. Avant toute mise en production, exécuter compute-sanitizer --tool memcheck et racecheck sur les cas de test, sur un jeu de données réduit.

Pièges courants

Le profiling n'est pas une étape ponctuelle en fin de développement : c'est un réflexe à réactiver à chaque changement significatif de kernel, de taille de problème ou de génération de GPU cible. Un kernel borné calcul sur une architecture peut devenir borné mémoire sur la suivante si le ratio entre débit de calcul et bande passante a changé — ce qui est systématiquement le cas d'une génération de carte à l'autre.