💻 computer science

Non-Uniform L2 Cache Latency Across the Streaming Multiprocessors of an NVIDIA L40

Cet article révèle que la latence de succès du cache L2 sur les GPU NVIDIA L40 et Blackwell est non uniforme et fortement dépendante du processeur de flux (SM) physique spécifique émettant la charge, permettant ainsi un primitif stable au niveau utilisateur pour l'autolocalisation des noyaux, l'empreinte digitale des périphériques et la distribution optimisée du travail, ce qui réduit le temps d'exécution (makespan) jusqu'à 11 %.

Auteurs originaux : Faruk Alpay, Baris Basaran

Publié 2026-06-23

📖 6 min de lecture🧠 Analyse approfondie

CC BY 4.0

Auteurs originaux : Faruk Alpay, Baris Basaran

Article original sous licence CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ Ceci est une explication générée par l'IA de l'article ci-dessous. Elle n'a pas été rédigée ni approuvée par les auteurs. Pour une précision technique, consultez l'article original. Lire la clause de non-responsabilité complète

Imaginez une bibliothèque massive et ultra-rapide (le GPU) où des milliers de bibliothécaires (appelés SM, ou processeurs de streaming multiprocessus) travaillent ensemble pour aller chercher des livres (des données) dans une immense salle de stockage partagée (le cache L2).

Pendant des années, les informaticiens ont cru que cette bibliothèque fonctionnait comme un système parfaitement uniforme : peu importe le bibliothécaire que vous sollicitiez, le temps nécessaire pour récupérer un livre dans la salle de stockage était exactement le même. Ils pensaient que la salle de stockage était un bassin d'informations unique et plat où la distance n'avait aucune importance.

Ce document révèle cependant que cette croyance est erronée. Les auteurs ont découvert que sur le GPU NVIDIA L40, l'endroit où se trouve physiquement un bibliothécaire détermine la vitesse à laquelle il peut récupérer un livre.

Voici le détail de leurs découvertes en utilisant des analogies simples :

1. La découverte du « l'importance de la distance »

Imaginez que la salle de stockage ne soit pas un sol plat, mais un grand entrepôt avec des allées.

L'ancienne vision : Tout le monde pensait que chaque bibliothécaire se trouvait exactement à la même distance des étagères. Si vous demandiez au Bibliothécaire A ou au Bibliothécaire Z, le temps d'attente était identique (environ 279 « ticks » de l'horloge).
La nouvelle réalité : Les auteurs ont mesuré le temps d'attente pour chacun des 142 bibliothécaires. Ils ont découvert que certains bibliothécaires se trouvent juste à côté des étagères et récupèrent les livres en 222 ticks. D'autres se trouvent à l'autre extrémité de l'entrepôt et mettent 339 ticks.
Le résultat : Cela représente une différence de vitesse de 52 %. C'est comme si une personne courait vers la porte d'entrée pendant qu'une autre doit courir jusqu'à la sortie arrière pour récupérer le même colis.

2. Le motif de l'« image miroir »

Lorsque les auteurs ont cartographié qui était rapide et qui était lent, ils n'ont pas vu un bruit aléatoire. Ils ont vu un motif parfait.

Les 142 bibliothécaires sont divisés en deux moitiés identiques (comme deux ailes d'un bâtiment).
Le Bibliothécaire n°1 de la première aile possède exactement le même profil de vitesse que le Bibliothécaire n°1 de la seconde aile.
Cette « symétrie en miroir » correspond au plan physique réel de la puce informatique, prouvant qu'il ne s'agit pas d'un bug logiciel — c'est un fait de la construction physique du matériel.

3. L'effet « empreinte digitale »

Parce que chaque bibliothécaire possède une vitesse unique basée sur son emplacement physique, la bibliothèque possède une identité secrète.

Auto-localisation : Si vous êtes un bibliothécaire (un programme informatique) et que vous demandez : « Combien de temps me faut-il pour récupérer un livre ? », vous pouvez instantanément déterminer l'endroit exact où vous vous trouvez dans l'entrepôt. Les auteurs ont construit un outil capable d'identifier votre emplacement spécifique avec une précision de 99 %.
Empreinte de l'appareil (Device Fingerprinting) : Même si vous possédez deux GPU L40 identiques et neufs (deux bibliothèques identiques), ils sont légèrement différents. Une bibliothèque peut avoir le Bibliothécaire n°5 placé légèrement plus près des étagères qu'une autre bibliothèque pour le même Bibliothécaire n°5, en raison de minuscules différences de fabrication. Les auteurs ont pu distinguer les deux machines identiques à 100 % en mesurant simplement ces infimes différences de vitesse. C'est comme être capable de distinguer des jumeaux identiques par leur façon de marcher.

4. Est-ce un risque de sécurité ?

Les auteurs prennent soin de préciser ce que cela implique pour la sécurité.

Ce que C'EST : Un moyen pour un programme de savoir où il se trouve à l'intérieur de l'ordinateur. C'est comme une personne entrant dans une pièce et réalisant : « Oh, je me tiens dans le coin près de la fenêtre ».
Ce que ce N'EST PAS : Ce n'est pas un moyen de voler les secrets d'autres programmes. Les auteurs soulignent que cela ne mesure que la propre vitesse du programme. Il ne peut pas espionner ce que font les autres programmes ni voler leurs données. C'est un outil d'« auto-localisation », pas un outil d'« espionnage ».

5. Le bénéfice pratique : Une planification plus intelligente

Pourquoi cela est-il important pour la performance ?

Imaginez que vous ayez une liste de 100 tâches à accomplir.
L'ancienne méthode : Vous assignez les tâches de manière aléatoire. Certaines vont à des bibliothécaires situés loin (lents), et d'autres à ceux qui sont proches (rapides). L'ensemble du travail attend que les personnes les plus lentes aient terminé.
La nouvelle méthode : Maintenant que nous avons la carte, nous pouvons confier le gros du travail aux bibliothécaires qui se trouvent les plus près des étagères.
Le résultat : En procédant ainsi, les auteurs ont montré qu'ils pouvaient terminer le travail 11 % plus rapidement pour les tâches qui dépendent fortement du cache. Cependant, si la tâche nécessite de récupérer des données dans la mémoire principale (un stockage différent et plus lent), cette astuce n'aide pas.

6. Ce n'est pas limité à une seule puce

Les auteurs ont testé cela sur une puce plus récente et différente (la RTX 5090) également. Ils ont constaté que la règle du « la distance compte » s'applique aussi là, bien que le motif soit légèrement différent. Cela prouve que l'idée ancienne d'un cache « uniforme » est probablement fausse pour beaucoup d'ordinateurs haut de gamme modernes.

Résumé

Le document brise l'illusion selon laquelle toutes les parties du cache d'un GPU sont égales. Il révèle que la localisation physique dicte la vitesse. En cartographiant ces différences de vitesse cachées, nous pouvons :

Identifier exactement où un programme est en train de s'exécuter.
Distinguer deux machines identiques.
Accélérer des tâches spécifiques en les assignant aux emplacements physiques les plus rapides.

Il s'avère que la puce informatique n'est pas un champ plat et uniforme ; c'est un paysage avec des collines et des vallées, et connaître la carte permet d'être plus rapide.

Résumé Technique : Latence de Cache L2 Non Uniforme sur NVIDIA L40

Énoncé du Problème
Le GPU NVIDIA L40, basé sur l'implémentation complète d'AD102, expose un cache L2 de 96 MiB qui est conventionnellement modélisé dans les outils de performance et la littérature académique comme un pool uniforme avec une latence de succès (hit latency) unique pour l'ensemble du dispositif. Cette abstraction suppose que tout multiprocesseur de flux (SM) accédant au cache L2 subit le même coût. Cependant, la réalité physique du cache L2 implique plusieurs tranches (slices) attachées à des contrôleurs mémoire, accédées via un réseau sur puce (NoC). Bien que des travaux antérieurs sur les architectures V100, A100 et H100 aient établi que la latence du NoC est non uniforme, ces études soulignaient que la latence moyenne par cœur était approximativement uniforme, la non-uniformité apparaissant principalement sous forme d'étapes discrètes à travers les frontières de partitions. Il restait une question ouverte de savoir si un GPU de centre de données à puce unique et partition unique de la génération Ada Lovelace (comme le L40) présente une structure comparable, et si, le cas échéant, la latence moyenne par SM est uniforme ou dépend du placement.

Méthodologie
Les auteurs ont développé une méthodologie de micro-benchmarking au niveau utilisateur pour mesurer la latence de succès L2 depuis chaque SM physique sans nécessiter de privilèges élevés ou de compteurs de performance matériels.

Conception de la sonde : La mesure utilise une chaîne de chargement dépendante à un seul thread et un seul warp (ld.global.cg) pour garantir un seul aller-retour à travers la structure du cache. Deux formes de tampons ont été employées : une forme à « ligne unique » pour sonder des tranches spécifiques de 128 B et une forme en « chaîne » (permutation aléatoire de 2 MiB) pour mesurer la distance moyenne vers toutes les tranches.
Isolation des SM : Pour éliminer la contention sur la puce, les blocs sont sérialisés à l'aide d'un compteur de tour global. Chaque bloc revendique un tour via un atomicAdd, attend, effectue une phase de chauffe (warm-up), puis mesure ses chargements dépendants de manière isolée.
Couverture : La sonde couvre l'ensemble des 142 SM physiques du L40 et les 256 sondes de tranches. Une sonde similaire a été exécutée sur un NVIDIA RTX 5090 (Blackwell GB202) pour tester la généralisation.
Modélisation : Les données de latence sont analysées à l'aide d'un modèle d'accès au cache non uniforme (NUCA) additif : $L = \mu + a(sm) + b(slice)$ , où $a(sm)$ représente le terme de placement du SM et $b(slice)$ représente le terme de placement de la tranche.

Résultats Clés

Variation Significative de la Latence : Contrairement au modèle uniforme, la latence de succès L2 sur le L40 n'est pas constante. Elle s'étend de 222,5 à 339,2 cycles (une plage de 52 %), avec une variation moyenne de 57,2 cycles entre les SM. Le plancher de bruit est négligeable (< 0,01 cycle), confirmant que la variation est physique.
Symétrie Bilatérale : Le terme de placement du SM présente une forte symétrie à deux faces. La division des 142 SM à l'identifiant 72 révèle deux moitiés (72 et 70 SM) ayant des profils de latence presque identiques (corrélation $r=0,999$ ). Cela concorde avec le plan de masse (floorplan) de l'AD102 composé de deux clusters GPC symétriques.
Gradient Continu vs Étapes Discrètes : Contra�à l'A100/H100, qui présentent une étape de latence discrète à la frontière de partition, le L40 présente un gradient de latence continu à travers les SM, sans écart bimodal.
Ajustement du Modèle Additif : Le modèle additif explique 87 % de la variation de latence ( $R^2=0,87$ ). L'ajout d'un seul terme d'interaction de rang 1 porte ce chiffre à 0,98, indiquant que la carte de latence est une géométrie de faible dimension suivant la disposition des GPC.
Généralisation et Spécificité :
- Trans-architecture : Le RTX 5090 montre également une non-uniformité (dispersion de 46 %), mais avec une symétrie plus faible ( $r=0,80$ ) et une latence absolue plus élevée.
- Empreinte Matérielle (Fingerprinting) : Deux L40 physiquement distincts, de modèle identique, sont distinguables avec une précision de 100 % sur la base de leurs cartes par SM, malgré des latences moyennes presque identiques (différence de 0,28 cycle). La carte par SM agit comme une identité matérielle stable.
Stabilité : La carte de latence par SM reste inchangée après une heure d'utilisation à 100 % sur les deux dispositifs, confirmant que la topologie est une propriété du silicium plutôt qu'un état transitoire.

Contributions et Signification
Le document revendique les contributions suivantes :

Cartographie de la Topologie : Il fournit la première carte de latence L2 par SM pour le L40, révélant une dispersion de 52 % qui invalide l'hypothèse de cache uniforme pour les charges de travail limitées par la latence.
Primitive d'Auto-Localisation : Les auteurs démontrent qu'une sonde de latence L2 unique permet à un noyau (kernel) d'identifier son propre placement physique avec une précision de 99,2 % (ou 92 % sur un ensemble de dispositifs mixtes) et d'identifier le die spécifique du dispositif avec une précision de 100 %. Cela est présenté comme une primitive d'« auto-localisation », et non comme une attaque contre un noyau victime, car il n'extrait aucune donnée secrète et repose sur le propre placement du noyau.
Implications de Performance : En distribuant le travail limité par la latence selon la carte par SM (en plaçant le travail sur les SM les plus rapides), les auteurs ont obtenu une réduction de 10,9 % du temps d'exécution (makespan) pour les charges de travail résidant dans le L2. Ce bénéfice disparaît dans les régimes limités par la DRAM, soulignant la nature spécifique au régime de l'impact de la topologie.
Réinterprétation de la Résidence : Le papier réinterprète les contrôles existants de résidence à thread unique (limites de capacité, mise à l'échelle de l'étiquette de ligne et résultats de persistance L2) comme des moyennes sur cette distribution non uniforme, plutôt que comme des constantes.

Portée et Limites
Les auteurs limitent explicitement la portée de leurs affirmations de sécurité. La technique permet l'auto-localisation (un noyau lisant son propre placement) et l'empreinte de l'appareil (fingerprinting). Elle ne permet pas d'inférer le placement d'un noyau « victime » co-résident, ni d'extraire des données secrètes ; de telles capacités nécessiteraient un canal de contention ou d'interférence qui n'est pas mesuré dans cette étude. L'étude est limitée à deux L40 et un RTX 5090, et la correspondance entre le %smid et le plan de masse physique est déduite de la symétrie et du nombre de GPC plutôt que confirmée indépendamment par photographie de die. Les comptes de cycles absolus sont spécifiques au micro-benchmark conçu et ne se traduisent pas directement en débit d'application où des requêtes indépendantes et l'usage de la L1 interviennent.