Aller au contenu

4 · Profilage avec Nsight

Deux outils, deux échelles. Nsight Systems répond à « où va le temps de mon application ». Nsight Compute répond à « pourquoi ce noyau est lent ». Les utiliser dans le mauvais ordre fait perdre des journées.

4.1 L'ordre correct

1. PyTorch profiler / chronomètre   → quel morceau du programme est lent ?
2. Nsight Systems                   → GPU inactif ? transferts ? synchronisation ?
3. Nsight Compute                   → ce noyau précis, pourquoi ?
4. SASS / cuobjdump                 → dernier recours

L'erreur qui coûte le plus cher

Ouvrir Nsight Compute en premier. C'est un microscope : il vous dira que votre noyau atteint 82 % de la bande passante — information exacte et inutile si le GPU est inactif 70 % du temps parce que le chargement des données est le goulot.

Toujours vérifier d'abord que le GPU travaille, avant de chercher à le faire travailler mieux.

4.2 Nsight Systems : la chronologie

nsys profile -o rapport --trace=cuda,nvtx,osrt,cudnn,cublas python train.py
nsys stats rapport.nsys-rep          # résumé texte

Puis ouverture de rapport.nsys-rep dans l'interface graphique.

Les quatre questions à poser à la chronologie

1. Y a-t-il des trous ?

GPU : [noyau][    trou 8ms    ][noyau][ trou ][noyau]
CPU : [........ Python ........][......][.....]

Un GPU inactif signifie que le CPU ne suit pas. Causes usuelles : chargement de données trop lent (augmentez num_workers), prétraitement en Python, synchronisations implicites.

2. Y a-t-il des synchronisations implicites ?

En PyTorch, ces appels forcent une synchronisation :

x.item()          # ← synchronise
x.cpu()           # ← synchronise
print(tenseur)    # ← synchronise
if tenseur > 0:   # ← synchronise
loss.item()       # ← le classique, dans la boucle d'entraînement

Chacun vide le pipeline. Une seule ligne print(loss.item()) par itération peut coûter 20 % du temps d'entraînement.

3. Les transferts sont-ils recouverts ?

Cherchez les barres MemcpyHtoD / MemcpyDtoH. Si elles alternent avec les noyaux au lieu de se superposer, il manque des flux ou de la mémoire épinglée.

4. Quels noyaux dominent ?

nsys stats --report cuda_gpu_kern_sum rapport.nsys-rep

 Time(%)  Total Time   Instances   Avg (ns)   Name
 --------  -----------  ---------  ---------  -----------------------
    38.2    1,204,331        512    2,352.2   ampere_bf16_gemm...
    17.9      564,120      1,024      550.9   elementwise_kernel...
    12.1      381,442        512      745.0   softmax_warp_forward...

C'est cette liste qui décide de ce que vous ouvrirez dans Nsight Compute.

Annoter avec NVTX

import torch.cuda.nvtx as nvtx

nvtx.range_push("forward")
sortie = modele(entree)
nvtx.range_pop()

nvtx.range_push("backward")
perte.backward()
nvtx.range_pop()

En C++ :

#include <nvtx3/nvToolsExt.h>
nvtxRangePushA("chargement des poids");
// ...
nvtxRangePop();

Sans NVTX, une chronologie de modèle réel est illisible : des milliers de noyaux aux noms générés. Avec, elle devient une carte.

4.3 Nsight Compute : le microscope

# Profil complet d'un noyau donné
ncu --set full -k mon_noyau -o rapport ./programme

# Cibler par nom (regex) et limiter le nombre d'instances
ncu --kernel-name regex:".*gemm.*" --launch-count 3 -o rapport ./programme

# Sections spécifiques (beaucoup plus rapide)
ncu --section SpeedOfLight --section MemoryWorkloadAnalysis ./programme

# Comparaison de deux versions
ncu --set full -o avant ./v1
ncu --set full -o apres ./v2
ncu-ui avant.ncu-rep apres.ncu-rep     # mode différentiel

Le profilage est intrusif

Nsight Compute rejoue chaque noyau plusieurs fois pour collecter les compteurs, et sérialise les lancements. Un programme profilé peut être 10 à 100× plus lent. Ce n'est pas un problème (les métriques restent justes), mais ne profilez pas une boucle d'entraînement entière : ciblez.

Il faut aussi les droits de lecture des compteurs : 

sudo sh -c 'echo "options nvidia NVreg_RestrictProfilingToAdminUsers=0" \
    > /etc/modprobe.d/nvidia-profiler.conf'

4.4 Lire un rapport, section par section

Speed of Light : le tableau de bord

C'est la première section, et souvent la seule nécessaire.

Section: GPU Speed Of Light Throughput
    Compute (SM) Throughput  [%]     23.14
    Memory Throughput        [%]     87.62
    DRAM Throughput          [%]     87.62
    L1/TEX Cache Throughput  [%]     41.20
    L2 Cache Throughput      [%]     55.31
    Duration                 [us]    142.30

Interprétation immédiate :

Compute Memory Diagnostic
bas haut limité par la mémoire — réduire le trafic, fusionner
haut bas limité par le calcul — précision réduite, tensor cores
bas bas limité par la latence — pas assez de parallélisme
haut haut équilibré, proche de l'optimum

L'exemple ci-dessus : 87,6 % de la mémoire, 23 % du calcul → limité par la mémoire, et déjà bien optimisé. Le seul levier restant est de transférer moins d'octets (fusion, quantification), pas d'optimiser le noyau.

Le cas « bas / bas » est le plus instructif

Si les deux sont bas, ni la mémoire ni le calcul ne sont saturés : le noyau attend. Causes : pas assez de blocs, trop de synchronisations, chaînes de dépendances longues, ou noyau trop court pour amortir son lancement.

C'est le régime dans lequel se trouvent la plupart des noyaux d'un décodage LLM à lot 1 — et c'est exactement ce que les megakernels corrigent.

Memory Workload Analysis

Le détail du trafic, niveau par niveau.

    Memory Throughput            [Gbyte/s]     2,935.21
    Mem Busy                     [%]              87.62
    Max Bandwidth                [%]              87.62
    L1/TEX Hit Rate              [%]              12.40
    L2 Hit Rate                  [%]              38.10
    Mem Pipes Busy               [%]              31.05

Métriques à surveiller :

Métrique Ce qu'elle révèle
dram__bytes.sum octets réellement lus/écrits en HBM — comparez à votre \(Q\) théorique
Secteurs par requête coalescence : 4 est parfait, 32 est catastrophique
l1tex__data_bank_conflicts_* conflits de banc en mémoire partagée
Local Memory non nul spilling de registres
L2 Hit Rate réutilisation inter-blocs

Le taux de réussite du L2 comme signal de localité de die

Sur les GPU multi-puces (B200, MI300X), un L2 Hit Rate faible avec des données qui devraient tenir en cache signale un problème de localité de chiplet : les blocs qui partagent des données sont dispersés sur des dies différents.

Fleet mesure exactement cela : en rendant l'ordonnancement conscient des chiplets, le taux passe de 12 % à 54 % (lot 32) et de 39 % à 61 % (lot 64), avec jusqu'à 37 % de trafic HBM en moins.

Warp State Statistics

Traité au chapitre précédent. Rappel du signal principal : Stall Long Scoreboard dominant = attente mémoire ; Stall Not Selected dominant = trop de warps.

Source Counters

Avec -lineinfo à la compilation, Nsight Compute attribue les compteurs ligne par ligne :

  Line  Instructions  Stall Long Scoreboard   Source
   142        45,120                  2,331   float a = A[i * N + k];
   143        45,120                 89,442   float b = B[k * N + j];   ← ici
   144        90,240                    112   acc += a * b;

La ligne 143 concentre les blocages : c'est l'accès non coalescé. C'est la fonctionnalité la plus utile de l'outil, et elle exige -lineinfo.

Occupancy

    Theoretical Occupancy        [%]     25.00
    Achieved Occupancy           [%]     23.14
    Block Limit Registers        [block]     4
    Block Limit Shared Mem       [block]     8
    Block Limit Warps            [block]     8

Les trois dernières lignes disent quelle ressource limite : ici les registres (4 blocs contre 8 pour les autres contraintes).

4.5 Les règles automatiques

Nsight Compute émet des diagnostics textuels. Ils sont bons et sous-utilisés :

[Warning] Uncoalesced Global Accesses
    The memory access pattern for global loads might not be optimal.
    On average, only 4.0 of the 32 bytes transmitted per sector are utilized.
    Est. Speedup: 62.4%

[Warning] Shared Memory Bank Conflicts
    The memory access pattern for shared loads causes 8-way bank conflicts.
    Est. Speedup: 21.1%

[Warning] Low Occupancy
    Achieved occupancy of 23.1% is below theoretical 25.0%.

Le champ Est. Speedup est une estimation de ce que vous gagneriez en corrigeant le point. Il permet de prioriser sans réfléchir. Commencez toujours par le plus élevé.

4.6 Profiler du PyTorch

Souvent le vrai point d'entrée.

import torch
from torch.profiler import profile, ProfilerActivity, schedule

with profile(
    activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
    schedule=schedule(wait=1, warmup=1, active=3),
    on_trace_ready=torch.profiler.tensorboard_trace_handler("./log"),
    record_shapes=True,
    with_stack=True,
) as prof:
    for i in range(5):
        sortie = modele(entree)
        perte = critere(sortie, cible)
        perte.backward()
        prof.step()

print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))

Ce que ça donne rapidement :

  • la répartition du temps entre opérateurs ;
  • les formes des tenseurs (record_shapes=True), indispensables pour savoir si une GEMM est bien dimensionnée ;
  • la trace exportable vers chrome://tracing ou TensorBoard ;
  • avec with_stack=True, la ligne Python responsable.

Le réflexe PyTorch

Avant tout, vérifiez trois choses :

  1. torch.backends.cuda.matmul.allow_tf32 = True (si la précision le permet) ;
  2. torch.compile(modele) — souvent 1,3 à 2× gratuitement grâce à la fusion ;
  3. aucun .item() ni print() dans la boucle chaude.

Ces trois points donnent en général plus que des jours de noyaux manuels.

4.7 Le SASS, en dernier recours

ncu --set full --import-source yes -o rapport ./prog
# puis dans l'interface : onglet "Source", vue "SASS"

ou hors ligne :

cuobjdump -sass ./binaire > sortie.sass
nvdisasm -c -g noyau.cubin

Ce qu'on y cherche concrètement :

Recherche Instruction SASS
Les tensor cores sont-ils utilisés ? HMMA, QGMMA, UTCHMMA
Les copies asynchrones ? LDGSTS (Ampere), UTMALDG (TMA Hopper)
Y a-t-il du spilling ? LDL / STL (local load/store)
Les accès sont-ils vectorisés ? LDG.E.128 contre LDG.E.32
Le déroulage a-t-il eu lieu ? compter les FFMA entre deux branchements

C'est un savoir-faire de spécialiste. Mais une seule recherche de HMMA dans le SASS répond en dix secondes à la question « mon noyau utilise-t-il vraiment les tensor cores ? », qui est autrement difficile à trancher.

4.8 Mesurer sans profileur, correctement

Parfois il faut juste un chiffre.

template <typename F>
float mesurer_us(F&& f, int repetitions = 100, int echauffement = 10) {
    for (int i = 0; i < echauffement; ++i) f();
    cudaDeviceSynchronize();

    cudaEvent_t d, fin;
    cudaEventCreate(&d); cudaEventCreate(&fin);
    cudaEventRecord(d);
    for (int i = 0; i < repetitions; ++i) f();
    cudaEventRecord(fin);
    cudaEventSynchronize(fin);

    float ms; cudaEventElapsedTime(&ms, d, fin);
    cudaEventDestroy(d); cudaEventDestroy(fin);
    return ms * 1000.0f / repetitions;
}

Les cinq règles d'une mesure valide :

  1. échauffer (JIT, allocations, fréquence) ;
  2. répéter et moyenner ;
  3. synchroniser avant de lire l'horloge ;
  4. verrouiller les fréquences (nvidia-smi -lgc) pour la reproductibilité ;
  5. vider les caches entre répétitions si le noyau est censé lire depuis la HBM — sinon vous mesurez le L2.

Le point 5 est souvent oublié et fausse tout : un noyau qui lit 20 Mo semble atteindre 10 To/s parce que les données sont restées en L2 (50 Mo sur H100).

Résumé du chapitre

À retenir

  • Ordre obligatoire : chronomètre → Nsight Systems → Nsight Compute → SASS.
  • Nsight Systems répond à « le GPU travaille-t-il ? ». Cherchez les trous, les synchronisations implicites, les transferts non recouverts.
  • Dans Nsight Compute, Speed of Light suffit dans 80 % des cas : compute contre memory, et le cas « bas/bas » qui signale une limitation par la latence.
  • Compilez avec -lineinfo : l'attribution ligne par ligne est la fonctionnalité la plus utile de l'outil.
  • Les règles automatiques donnent un Est. Speedup : traitez-les par ordre décroissant.
  • En PyTorch : TF32, torch.compile, et pas de .item() dans la boucle, avant toute optimisation manuelle.
  • Mesurer, c'est : échauffer, répéter, synchroniser, verrouiller les fréquences, vider les caches.

Vérifiez que vous avez compris

Speed of Light indique Compute 8 %, Memory 11 %. Que faire ?

Ni l'un ni l'autre n'est saturé : le noyau est limité par la latence ou par le lancement.

Vérifiez, dans l'ordre :

  1. la durée du noyau. Si elle est de quelques microsecondes, le surcoût de lancement domine → fusionner avec les noyaux voisins ;
  2. le nombre de blocs. S'il est inférieur au nombre de SM, une partie du GPU ne fait rien → augmenter le parallélisme ;
  3. les états de warp. Stall Barrier élevé → trop de __syncthreads() ou déséquilibre.

C'est le profil typique des petits noyaux élémentaires d'un modèle, et la justification directe de la fusion puis des megakernels.

Votre noyau atteint 2 900 Go/s sur H100 mais le calcul théorique dit qu'il ne devrait transférer que 40 Mo, pour un temps minimal de 12 µs. Il en met 45. Où est le problème ?

2 900 Go/s représente 87 % de la bande passante : le noyau sature la mémoire. Mais s'il met 45 µs, c'est qu'il transfère \(2\,900 \times 10^9 \times 45 \times 10^{-6} = 130\) Mo, soit 3,3× plus que les 40 Mo utiles.

Diagnostic : accès non coalescés (facteur ~3-4 est cohérent avec un pas de 2 ou un désalignement), ou données relues plusieurs fois faute de réutilisation en mémoire partagée.

Vérification : comparez dram__bytes.sum à 40 Mo. Le rapport donne directement \(1/\eta\).

Pourquoi --set full peut-il rendre un programme 100× plus lent, alors que les résultats restent justes ?

Nsight Compute collecte des centaines de compteurs matériels, dont beaucoup ne peuvent pas être lus simultanément. Il rejoue donc chaque noyau plusieurs fois (replay), en réinitialisant l'état entre les passes, et sérialise tous les lancements pour isoler les mesures.

Les métriques restent justes parce que chaque passe mesure le noyau dans des conditions contrôlées. Ce qui n'est plus juste, c'est la durée totale du programme et tout ce qui dépend du recouvrement entre noyaux — d'où l'utilité de Nsight Systems, qui lui n'instrumente que les frontières.

Chapitre suivant : 5 · Checklist d'optimisation

Sources de ce chapitre