Aller au contenu

6 · Fleet et les chiplets

Depuis 2024, « un GPU » n'est plus une puce. Le B200 en compte deux, le MI300X huit. Le modèle de programmation, lui, n'a pas bougé — et c'est un problème.

6.1 Le constat

Le résumé de Fleet le formule directement :

Les GPU modernes adoptent des conceptions à base de chiplets avec plusieurs hiérarchies de caches privées, mais les modèles de programmation actuels (CUDA/HIP) exposent une hiérarchie d'exécution plate qui ne peut pas exprimer la localité ni la synchronisation au niveau du chiplet.

Concrètement :

GPU Structure physique Ce que CUDA/HIP expose
H100 1 die, L2 partagé (2 partitions) 132 SM, un L2
B200 2 dies liés par NV-HBI, L2 en 4 partitions 148 SM, « un » L2
MI300X 8 XCD, chacun avec son L2 304 CU, « un » L2
MI350X 8 XCD, 256 CU idem

Le programmeur ne peut pas dire « place ces deux blocs sur le même die », ni « synchronise seulement les blocs de mon chiplet ».

La conséquence : deux blocs qui partagent des données peuvent atterrir sur des dies différents, et chaque accès partagé devient un défaut de L2 suivi d'un accès HBM.

6.2 Pourquoi cela touche particulièrement les megakernels

Un megakernel occupe tout le GPU avec des blocs persistants qui communiquent intensivement entre eux. C'est le pire cas possible pour un problème de localité de die :

  • les activations produites par un bloc sont consommées par d'autres ;
  • les compteurs de dépendance sont sondés depuis partout ;
  • les poids d'une couche sont lus par de nombreux blocs.

Dans un modèle d'exécution plat, ces communications traversent les frontières de die de façon arbitraire.

6.3 Le papier

Titre : Fleet: Hierarchical Task-based Abstraction for Megakernels on Multi-Die GPUs

Date : 20 avril 2026 (arXiv:2604.15379)

Auteurs : Sangeeta Chowdhary, Ryan Swann, Sean Siddens, Muhammad Osama, Stephen Neuendorffer, Alexandru Dutu, Karthik Sangaiah, Sandeepa Bhuyan, Samuel Bayliss, Ganesh Dasika

Les affiliations sont principalement AMD, ce qui est cohérent avec la cible (MI350) et avec le fait que le problème des chiplets y est le plus aigu.

6.4 L'abstraction : les chiplet-tasks

Fleet introduit une notion de tâche liée à un chiplet (Chiplet-task), qui lie explicitement le travail et les données à un chiplet donné.

Modèle plat (CUDA/HIP)          Modèle Fleet
──────────────────────          ────────────
grille                          grille
  └─ bloc  (SM/CU arbitraire)     └─ chiplet-task  (lié à un XCD)
                                       └─ bloc

Le runtime est un noyau persistant avec un ordonnancement par chiplet, ce qui permet une exécution coopérative des tâches avec une réutilisation de cache coordonnée.

L'idée pratique : les tâches qui partagent des données sont assignées au même chiplet, de sorte que ces données restent dans le L2 local.

Cela peut impliquer de dupliquer certaines données par chiplet — le même arbitrage espace/temps que celui du monokernel de Kog.

6.5 Les résultats

Matériel : AMD Instinct MI350, modèle Qwen3-8B.

Métrique Résultat
Latence de décodage, lots 1-8 1,3 – 1,5× plus faible que vLLM
Taux de réussite du L2, lot 32 12 % → 54 %
Taux de réussite du L2, lot 64 39 % → 61 %
Réduction du trafic HBM jusqu'à 37 %
Contre un megakernel non conscient des chiplets, lots élevés 1,27 – 1,30×

Le chiffre le plus parlant

Le taux de réussite du L2 passe de 12 % à 54 % à lot 32.

Cela signifie qu'avant, 88 % des accès manquaient le L2 et allaient chercher en HBM — alors que les données étaient dans un L2, juste pas dans le bon.

C'est exactement le symptôme décrit en Performance 4 : un taux de réussite L2 anormalement bas avec des données qui devraient tenir en cache signale un problème de localité de die, invisible dans les métriques agrégées.

Et la ligne « 1,27-1,30× contre un megakernel non conscient des chiplets » est la plus importante méthodologiquement : elle isole le gain dû à la conscience des chiplets, indépendamment du gain dû au megakernel lui-même.

6.6 Le cas NVIDIA

Le problème existe aussi sur Blackwell, en moins prononcé.

Un B200 a deux dies reliés par NV-HBI à 10 To/s, présentés comme un seul GPU CUDA. Le microbenchmarking de Blackwell note que le L2 y est découpé en 4 partitions, contre 2 sur Hopper.

Deux différences avec AMD atténuent le problème :

  1. deux dies au lieu de huit : la probabilité qu'un partage traverse une frontière est plus faible ;
  2. NV-HBI à 10 To/s est bien plus rapide que l'interconnexion inter-XCD d'AMD.

Mais le problème est structurellement le même, et il s'aggravera : la trajectoire industrielle va vers davantage de chiplets, pas moins.

6.7 Ce qu'on peut faire aujourd'hui, sans Fleet

Trois techniques accessibles en CUDA/HIP standard.

1. Exploiter le mapping implicite des identifiants de bloc

Sur MI300X, les blocs sont distribués sur les XCD selon un motif en tourniquet sur blockIdx. On peut donc déduire le XCD d'un bloc :

int xcd = blockIdx.x % NB_XCD;      // vrai sur MI300X, à vérifier par carte

et organiser le travail en conséquence — placer les blocs qui partagent des données à des indices congruents modulo le nombre de XCD.

C'est un détail d'implémentation, pas un contrat

Ce mapping n'est pas documenté comme garanti. Il peut changer avec le pilote, la version de ROCm, ou la carte. Un code qui en dépend doit le vérifier au démarrage et prévoir un chemin de repli.

C'est précisément le genre de fragilité que Fleet cherche à remplacer par une abstraction propre.

2. Dupliquer les données partagées par die

C'est ce que fait Kog : « exploitant les 8 XCD du MI300X avec leurs propres caches L2, ils ont dupliqué les tenseurs par die d'E/S pour éviter les pénalités de communication inter-chiplets ».

Coût : \(N\) fois la mémoire pour les données concernées. Gain : accès L2 local au lieu de HBM.

3. Utiliser les clusters sur NVIDIA

Sur Hopper et Blackwell, les blocs d'un cluster sont garantis sur des SM du même GPC, donc du même die. Un cluster est donc un moyen indirect d'exprimer la localité de die.

C'est partiel — il n'y a pas de contrôle sur quel die — mais c'est une garantie de co-localisation.

6.8 Ce que cela dit du modèle de programmation

Une abstraction qui vieillit

Le modèle grille/bloc/thread date de 2006 et n'a été étendu qu'une fois, avec les clusters de Hopper (2022).

Pendant ce temps, le matériel a acquis : des unités matricielles, une mémoire de tenseurs, des moteurs de copie asynchrones, plusieurs dies, plusieurs partitions de cache, et des files de travail matérielles.

Fleet est un symptôme : le modèle plat ne suffit plus à exprimer ce que le matériel peut faire. C'est le même diagnostic que celui qui a produit CUDA Tile, Triton et les autres DSL — sous un angle différent.

La question ouverte : le prochain niveau de hiérarchie sera-t-il exposé par CUDA/HIP directement, ou seulement par des abstractions de plus haut niveau ?

Résumé du chapitre

À retenir

  • Un B200 a 2 dies, un MI300X 8 XCD, chacun avec son L2. CUDA et HIP exposent une hiérarchie plate qui ne peut pas exprimer cette localité.
  • Les megakernels sont le pire cas : blocs persistants communiquant intensivement à travers tout le GPU.
  • Fleet (AMD, avril 2026) introduit les chiplet-tasks, qui lient travail et données à un chiplet, avec un runtime persistant à ordonnancement par chiplet.
  • Résultats sur MI350/Qwen3-8B : 1,3-1,5× vs vLLM (lots 1-8), taux de réussite L2 de 12 % à 54 %, −37 % de trafic HBM, et 1,27-1,30× contre un megakernel non conscient des chiplets.
  • Sans Fleet : exploiter le mapping implicite blockIdx % NB_XCD (fragile), dupliquer les données par die, ou utiliser les clusters sur NVIDIA.
  • Le modèle grille/bloc/thread montre ses limites face au matériel de 2026.

Vérifiez que vous avez compris

Pourquoi le gain de Fleet est-il plus élevé aux petits lots contre vLLM, mais plus faible contre un megakernel non conscient des chiplets ?

Parce que les deux comparaisons mesurent des choses différentes.

Contre vLLM (1,3-1,5× aux lots 1-8) : le gain cumule deux effets — celui du megakernel lui-même (suppression des frontières, décisif à petit lot) et celui de la conscience des chiplets.

Contre un megakernel non conscient des chiplets (1,27-1,30× aux lots plus grands) : le gain isole le second effet. Il apparaît surtout aux lots élevés, où le volume de données partagées est important et où la localité de cache compte le plus.

Publier les deux comparaisons est méthodologiquement rigoureux : cela sépare la contribution propre du papier de celle de la technique de base.

Dupliquer un tenseur sur 8 XCD multiplie par 8 son empreinte mémoire. Quand est-ce rentable ?

Quand le tenseur est petit et très relu.

Le calcul : soit \(s\) la taille du tenseur et \(r\) le nombre de fois qu'il est lu par XCD.

  • Sans duplication : \(s\) octets stockés, mais \(7/8\) des lectures manquent le L2 local → environ \(7r s/8\) octets de trafic HBM.
  • Avec duplication : \(8s\) octets stockés, une lecture HBM par XCD pour remplir, puis tout en L2 local → \(8s\) octets de trafic HBM.

C'est rentable dès que \(7rs/8 > 8s\), soit \(r > 9{,}1\).

Un tenseur relu plus de ~9 fois par XCD mérite d'être dupliqué, à condition que \(8s\) tienne en mémoire. C'est typiquement le cas des poids d'une couche dans un megakernel, ou des tampons d'activation partagés.

Le mapping blockIdx % NB_XCD n'est pas documenté. Comment le vérifier ?

En le mesurant au démarrage, avec un micro-benchmark :

  1. lancer un noyau où chaque bloc écrit son identifiant matériel — sur AMD, via s_getreg_b32 pour lire le registre matériel identifiant le CU et le XCD ;
  2. construire la table blockIdx → XCD ;
  3. vérifier qu'elle correspond au motif attendu ;
  4. si oui, activer le chemin optimisé ; sinon, utiliser un chemin générique.

C'est une calibration à l'exécution, comparable à ce qu'on fait pour tout paramètre matériel non contractuel. Le coût est de quelques microsecondes au démarrage, et cela évite qu'une mise à jour de pilote ne dégrade silencieusement les performances.

Chapitre suivant : 7 · AMD et le monokernel

Sources de ce chapitre