Aller au contenu

3 · Occupancy et latence

Le mythe le plus tenace de la programmation GPU : « il faut maximiser l'occupancy ». C'est faux depuis 2010, et la démonstration est instructive.

3.1 Définition

L'occupancy est le rapport entre le nombre de warps résidents sur un SM et le maximum matériel :

\[ \text{occupancy} = \frac{\text{warps résidents}}{\text{warps max par SM}} \]

Sur Hopper et Blackwell, le maximum est 64 warps (2 048 threads) par SM.

Trois ressources la limitent :

Ressource Contrainte
Registres \(\text{warps} \le \lfloor 65\,536 / (32 R) \rfloor\), avec \(R\) registres/thread
Mémoire partagée \(\text{blocs} \le \lfloor S_{\text{SM}} / S_{\text{bloc}} \rfloor\)
Blocs \(\le 32\) blocs résidents par SM (Hopper)

Le minimum des trois gagne.

3.2 Pourquoi on croyait qu'il en fallait beaucoup

L'argument classique, correct dans son principe : l'occupancy sert à cacher la latence.

Formalisons avec la loi de Little. Pour saturer un système de latence \(L\) et de débit \(T\), il faut maintenir en vol :

\[ N = L \times T \quad \text{opérations concurrentes} \]

Appliquons-la à la mémoire d'un H100 :

  • latence mémoire \(L \approx 500\) cycles ;
  • bande passante \(B = 3{,}35\) To/s à ~1,7 GHz, soit ~1 970 octets par cycle pour tout le GPU, ou ~15 octets par cycle par SM.

Pour saturer la mémoire, il faut donc environ :

\[ 500 \times 15 = 7\,500 \text{ octets en vol par SM} \]

Si chaque thread a une lecture de 4 octets en vol, il faut 1 875 threads — proche du maximum de 2 048. D'où la conclusion : occupancy maximale.

Le raisonnement est juste, mais son hypothèse ne l'est pas.

3.3 La démonstration de Volkov

En 2010, Vasily Volkov (Berkeley) présente à la GTC un exposé intitulé Better Performance at Lower Occupancy. Son observation :

L'occupancy mesure le parallélisme au niveau des threads (TLP). Mais le parallélisme au niveau des instructions (ILP) cache la latence tout autant.

Si chaque thread a quatre lectures indépendantes en vol au lieu d'une, il faut quatre fois moins de threads pour maintenir le même nombre d'octets en vol.

\[ \text{octets en vol} = \text{threads} \times \text{ILP} \times \text{taille d'accès} \]

Ses chiffres, sur le matériel de l'époque : avec un ILP de 3, 256 threads par SM suffisent à atteindre 100 % d'utilisation ; avec un ILP de 4, 192 threads. Contre 1 536 dans le régime ILP = 1.

Et le gain est double, parce que baisser l'occupancy libère des registres :

\[ R_{\max} = \frac{65\,536}{32 \times \text{warps résidents}} \]
Warps résidents Occupancy Registres/thread disponibles
64 100 % 32
32 50 % 64
16 25 % 128
8 12,5 % 255 (limite matérielle)

Plus de registres = plus de valeurs gardées près du calcul = moins d'accès mémoire = moins de latence à cacher. Le cercle est vertueux.

L'énoncé correct

Il ne faut pas maximiser l'occupancy. Il faut maintenir assez d'accès mémoire en vol pour saturer la bande passante, ce qui s'obtient par un mélange de TLP et d'ILP. Quand l'ILP est élevé, une occupancy faible est non seulement acceptable mais préférable.

3.4 Comment créer de l'ILP

Trois techniques, toutes déjà rencontrées.

1. Plusieurs éléments par thread.

// ILP = 1
float x = a[i];
c[i] = x * 2.0f;

// ILP = 4 : quatre lectures indépendantes en vol simultanément
float x0 = a[i], x1 = a[i + s], x2 = a[i + 2*s], x3 = a[i + 3*s];
c[i] = x0 * 2.0f;  c[i + s] = x1 * 2.0f;
c[i + 2*s] = x2 * 2.0f;  c[i + 3*s] = x3 * 2.0f;

2. Le chargement vectorisé.

float4 v = reinterpret_cast<const float4*>(a)[i];   // 16 octets en 1 instruction

Une seule instruction met 16 octets en vol au lieu de 4.

3. Le déroulage de boucle.

#pragma unroll 4
for (int k = 0; k < K; ++k) { acc += a[k] * b[k]; }

Le compilateur émet quatre chargements avant la première accumulation, créant naturellement de l'ILP.

Le cas de la GEMM

Un noyau CUTLASS de GEMM tourne typiquement à 12 à 25 % d'occupancy. Il utilise 100 à 200 registres par thread pour garder un bloc d'accumulateurs \(8\times8\) ou plus, et son ILP est énorme (64 FMA indépendantes par itération du produit extérieur).

C'est le contre-exemple parfait au mythe de l'occupancy : le noyau le plus optimisé du domaine a l'occupancy la plus basse.

3.5 Quand l'occupancy compte vraiment

Trois situations où il en faut beaucoup.

1. Noyaux à accès mémoire irréguliers. Parcours de graphes, tables de hachage, tri par indirection : la latence varie beaucoup et l'ILP est difficile à créer. Il faut alors beaucoup de warps pour absorber la variance.

2. Noyaux à très faible ILP par nature. Une chaîne de dépendances séquentielle (par exemple un scan intra-thread) n'offre aucun ILP.

3. Noyaux avec branchement divergent. La divergence réduit le nombre de threads utiles par warp ; plus de warps compensent.

3.6 Le diagnostic correct

L'occupancy n'est pas la métrique à regarder. Voici les bonnes.

Les états de warp (Warp State Statistics)

Nsight Compute décompose, pour chaque cycle, pourquoi les warps ne progressent pas :

État (stall reason) Signification Remède
Stall Long Scoreboard attente d'un accès mémoire globale plus d'ILP, ou réduire le trafic
Stall Short Scoreboard attente de mémoire partagée conflits de banc ?
Stall Barrier attente à __syncthreads() déséquilibre entre warps
Stall Wait dépendance sur une instruction courte plus d'ILP
Stall Not Selected prêt mais un autre warp a été choisi ← c'est bon signe
Stall MIO Throttle file d'attente des unités mémoire pleine trop de requêtes
Stall Math Pipe Throttle unité de calcul saturée ← limité par le calcul

La métrique reine

Stall Not Selected élevé signifie que vous avez trop de warps, pas trop peu : le matériel a l'embarras du choix. Si en plus Stall Long Scoreboard est faible, votre occupancy est plus que suffisante et vous pourriez l'échanger contre des registres.

À l'inverse, Stall Long Scoreboard dominant avec une occupancy déjà à 100 % signifie que le problème est le volume de trafic mémoire, pas sa latence. Retour au roofline.

Le débit atteint

La seule métrique qui ne ment jamais : comparez les Go/s ou les TFLOPS obtenus aux valeurs crêtes. Si vous êtes à 90 %, l'occupancy est bonne par définition, quelle que soit sa valeur.

3.7 Une méthode de décision

1. Mesurer le débit atteint (Go/s ou TFLOPS)
   └─ > 80 % du plafond pertinent → arrêter, c'est bon

2. Regarder les états de warp
   ├─ "Long Scoreboard" dominant
   │   ├─ occupancy < 50 % → augmenter (moins de registres, blocs plus petits)
   │   └─ occupancy > 50 % → augmenter l'ILP, ou réduire le trafic
   ├─ "Barrier" dominant
   │   └─ déséquilibre de charge : revoir le découpage
   ├─ "Not Selected" dominant
   │   └─ trop de warps : baisser l'occupancy, prendre plus de registres
   └─ "Math Pipe Throttle" dominant
       └─ limité par le calcul : tensor cores, précision réduite

3. Ne jamais optimiser l'occupancy pour elle-même

3.8 Le cas extrême : les noyaux persistants

Un megakernel pousse la logique à son terme : un seul bloc par SM, résident pendant toute l'exécution du modèle.

L'occupancy est alors fixée par la taille du bloc. Mirage MPK utilise 128 threads par worker sur H100, soit 4 warps sur 64 possibles : 6,25 % d'occupancy.

Est-ce un problème ? Non, parce que :

  • l'ILP est maximal (pipelining logiciel explicite, TMA en vol permanent) ;
  • toute la mémoire partagée du SM est disponible pour un seul bloc, ce qui permet des tuiles maximales ;
  • l'objectif n'est pas de cacher la latence par le TLP mais de ne jamais laisser la mémoire inactive, ce qui se fait par le pipelining asynchrone.

Résultat mesuré par Hazy Research : 78 % de la bande passante mémoire d'un H100, contre ~50 % pour les systèmes à noyaux multiples — avec une occupancy dérisoire.

La conclusion du chapitre

L'occupancy est un moyen de cacher la latence, parmi d'autres. La programmation GPU moderne (Hopper, Blackwell) remplace progressivement le TLP par de l'asynchronisme explicite : TMA, mbarrier, spécialisation des warps.

Autrement dit : l'ère où l'occupancy était la bonne métrique est en train de se refermer.

Résumé du chapitre

À retenir

  • Occupancy = warps résidents / 64. Limitée par les registres, la mémoire partagée, ou le nombre de blocs.
  • Loi de Little : ce qu'il faut, ce sont des accès en vol, obtenus par TLP ou ILP.
  • Volkov (2010) : avec ILP = 4, 192 threads par SM suffisent. Baisser l'occupancy libère des registres, ce qui réduit le trafic.
  • Les noyaux GEMM les plus rapides tournent à 12-25 % d'occupancy.
  • Diagnostiquez avec les états de warp, pas avec l'occupancy. Stall Not Selected élevé = trop de warps.
  • Les megakernels tournent à ~6 % d'occupancy et atteignent 78 % de la bande passante.

Vérifiez que vous avez compris

Votre noyau passe de 100 % à 50 % d'occupancy après une optimisation, et devient 1,4× plus rapide. Explication ?

L'optimisation a probablement augmenté le travail par thread (plus d'éléments, boucle déroulée), ce qui a augmenté la consommation de registres et donc réduit l'occupancy — mais a créé de l'ILP et réduit le nombre total d'instructions.

C'est le scénario classique décrit par Volkov. L'occupancy a baissé parce que chaque thread fait plus de travail utile. Le seul indicateur qui compte est le 1,4×.

Nsight Compute indique Achieved Occupancy: 23% et Theoretical Occupancy: 25%. Bon ou mauvais ?

L'écart entre théorique et atteinte est faible (2 points), ce qui est bon : cela signifie que les warps sont effectivement résidents, sans déséquilibre de charge ni queue de vagues significative.

La valeur absolue de 25 % ne dit rien en soi. Elle indique juste que la consommation de registres ou de mémoire partagée limite à 16 warps. Pour savoir si c'est un problème, il faut le débit atteint et les états de warp.

Un écart important entre théorique et atteinte (par exemple 25 % contre 12 %) est en revanche un signal : quantification de vagues, déséquilibre, ou blocs qui se terminent à des instants très différents.

Pourquoi la spécialisation des warps (producteurs/consommateurs) réduit-elle le besoin d'occupancy ?

Parce qu'elle remplace le TLP par de l'asynchronisme explicite. Dans un schéma warp-specialized, un warp dédié émet en permanence des copies TMA dont les données arrivent pendant que d'autres warps calculent. Le nombre d'octets en vol est piloté par la profondeur du pipeline logiciel (nombre d'étages de tampon), pas par le nombre de warps.

On passe d'un régime « beaucoup de warps qui attendent chacun un peu » à « peu de warps, mais un flux continu de données organisé à la main ». C'est plus efficace et bien plus difficile à écrire — d'où les bibliothèques. Voir CUDA moderne · Spécialisation des warps.

Chapitre suivant : 4 · Profilage avec Nsight

Sources de ce chapitre