Aller au contenu

3 · Clusters et mémoire partagée distribuée

Hopper a ajouté un niveau à une hiérarchie qui n'avait pas bougé depuis 2006. Entre le bloc et la grille s'intercale désormais le cluster : un groupe de blocs garantis co-résidents, capables de se synchroniser et de lire la mémoire partagée les uns des autres.

3.1 La hiérarchie, mise à jour

grille
  └─ cluster            ← nouveau (Hopper, sm_90)
       └─ bloc (CTA)
            └─ warp
                 └─ thread

Un cluster (aussi appelé Cooperative Thread Array Cluster) :

  • regroupe de 1 à 8 blocs (16 sur certaines configurations non portables) ;
  • ses blocs sont garantis simultanément résidents sur des SM du même GPC (Graphics Processing Cluster) ;
  • ses blocs peuvent se synchroniser entre eux ;
  • ses blocs peuvent lire, écrire et faire des atomiques dans la mémoire partagée les uns des autres.

Cette dernière propriété s'appelle la mémoire partagée distribuée (DSMEM, Distributed Shared Memory). Elle repose sur un réseau d'interconnexion SM-à-SM ajouté dans Hopper.

3.2 Déclarer et lancer un cluster

Deux façons.

Statique, dans l'attribut du noyau :

__global__ void __cluster_dims__(2, 1, 1) mon_noyau(...) { ... }
mon_noyau<<<grille, bloc>>>(...);

Dynamique, à travers l'API de lancement étendue :

cudaLaunchConfig_t config = {0};
config.gridDim  = grille;
config.blockDim = bloc;

cudaLaunchAttribute attr[1];
attr[0].id = cudaLaunchAttributeClusterDimension;
attr[0].val.clusterDim.x = 2;
attr[0].val.clusterDim.y = 1;
attr[0].val.clusterDim.z = 1;
config.attrs    = attr;
config.numAttrs = 1;

cudaLaunchKernelEx(&config, mon_noyau, arg1, arg2);

La taille de grille doit être un multiple de la taille de cluster

Si gridDim.x = 100 et clusterDim.x = 3, le lancement échoue. Le cluster est une unité indivisible d'ordonnancement.

Un cluster de 8 est portable seulement s'il est déclaré comme tel ; au-delà, il faut vérifier cudaOccupancyMaxPotentialClusterSize.

3.3 Utiliser la mémoire partagée distribuée

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void __cluster_dims__(2, 1, 1) exemple() {
    __shared__ float tuile[1024];

    cg::cluster_group cluster = cg::this_cluster();
    unsigned mon_rang = cluster.block_rank();       // 0 ou 1
    unsigned voisin   = mon_rang ^ 1;

    // Remplir ma propre tuile
    tuile[threadIdx.x] = calculer(threadIdx.x);

    cluster.sync();          // ← barrière sur les 2 blocs du cluster

    // Obtenir un pointeur vers la mémoire partagée du bloc voisin
    float* tuile_voisine = cluster.map_shared_rank(tuile, voisin);

    // Lire directement dedans, sans passer par la mémoire globale
    float v = tuile_voisine[threadIdx.x];

    cluster.sync();          // avant que quiconque écrase sa tuile
}

C'est cela, l'apport fondamental : deux blocs échangent des données sans aller-retour en mémoire globale.

Chemin Latence approx. Bande passante
Mémoire partagée locale ~25 cycles très élevée
DSMEM (SM voisin, même GPC) ~50-70 cycles élevée
L2 ~200 cycles ~7 To/s
HBM ~500 cycles 3,35 To/s

DSMEM se situe entre la mémoire partagée locale et le L2 — bien plus près de la première.

3.4 À quoi ça sert concrètement

1. Agrandir la tuile effective

Un bloc Hopper dispose de 227 Ko de mémoire partagée. Un cluster de 4 blocs dispose de 908 Ko collectivement adressables.

Pour une GEMM, cela permet des tuiles quatre fois plus grandes sans changer la taille des blocs, donc une intensité arithmétique quatre fois supérieure.

2. Diviser le trafic HBM par le multicast TMA

Comme vu au chapitre précédent, une seule copie TMA peut alimenter tous les blocs d'un cluster. Sans clusters, chaque bloc lit sa propre copie depuis la HBM.

C'est probablement l'usage le plus rentable en pratique.

3. Réduire la profondeur des réductions

Une réduction sur une grille se fait classiquement en deux noyaux (partiels puis final). Avec des clusters, on peut réduire au niveau du cluster puis n'émettre qu'un atomique par cluster au lieu d'un par bloc.

4. Coopérer sur une MMA (Blackwell)

Sur Blackwell, tcgen05.mma en mode paire de CTA exige que les deux blocs soient dans le même cluster. La tuile passe de \(128\times128\) à \(256\times128\). Voir le chapitre suivant.

3.5 Un exemple : GEMM avec cluster

Le schéma, tel qu'utilisé par CUTLASS sur Hopper et Blackwell :

Cluster de 2 blocs, calculant deux tuiles de C côte à côte

           ┌─────────── A (partagé par les 2 blocs) ───────────┐
           │  chargé UNE fois par TMA multicast                │
           └──────────────────────────────────────────────────┘
                    │                        │
              ┌─────▼─────┐            ┌─────▼─────┐
              │  Bloc 0   │            │  Bloc 1   │
              │  B tuile0 │            │  B tuile1 │
              │  → C[0]   │            │  → C[1]   │
              └───────────┘            └───────────┘

Le trafic HBM pour \(\mathbf{A}\) est divisé par deux. Le tutoriel Colfax sur les GEMM avec clusters sur Blackwell détaille l'implémentation complète.

3.6 Les contraintes et les pièges

Ce qu'il faut savoir avant d'utiliser des clusters

1. Tous les blocs du cluster doivent atteindre cluster.sync(). Sinon, interblocage. Comme __syncthreads(), mais à une échelle où le débogage est plus pénible.

2. Le pointeur DSMEM n'est valide que pendant la vie du cluster. map_shared_rank ne doit jamais être stocké au-delà.

3. La taille du cluster réduit la flexibilité d'ordonnancement. Un cluster de 8 blocs ne peut démarrer que si 8 SM du même GPC sont libres simultanément. Sur un GPU chargé, cela peut allonger le temps d'attente.

4. Ce n'est pas portable. sm_90 minimum. Aucun équivalent AMD direct — c'est d'ailleurs une des raisons pour lesquelles Fleet a dû inventer les « chiplet-tasks » pour exprimer la localité sur MI350.

5. cluster.sync() n'est pas gratuit. De l'ordre de la centaine de cycles. Beaucoup moins qu'une barrière de grille (des microsecondes), beaucoup plus qu'un __syncthreads().

3.7 Cluster Launch Control (Blackwell)

Blackwell ajoute un mécanisme qui mérite d'être connu pour la partie 8 : le Cluster Launch Control.

Dans le modèle classique, un bloc connaît son travail par son blockIdx : le mapping est statique. Avec CLC, un cluster peut demander au matériel le prochain identifiant de travail disponible, ce qui permet un ordonnancement dynamique piloté par le GPU.

Modèle statique          Modèle CLC
───────────────          ──────────
bloc 17 → tuile 17       cluster demande → matériel répond "tuile 43"
                         cluster demande → matériel répond "tuile 51"
                         …

Pourquoi c'est important : dans un noyau persistant où les tâches ont des durées très inégales, l'affectation statique crée un déséquilibre. CLC fournit une file de travail matérielle, sans le coût des atomiques logiciels.

C'est exactement le problème que les megakernels résolvent en logiciel avec leurs files de tâches (Mirage MPK, la file globale de Hazy Research pour le megakernel de débit). CLC en fournit une version matérielle.

Résumé du chapitre

À retenir

  • Le cluster est un niveau entre le bloc et la grille : 1 à 8 blocs garantis co-résidents sur des SM du même GPC.
  • Ils peuvent se synchroniser (cluster.sync(), ~100 cycles) et lire la mémoire partagée les uns des autres (DSMEM, ~50-70 cycles).
  • Trois usages rentables : agrandir la tuile effective, diviser le trafic HBM par le multicast TMA, et permettre les MMA sur paire de CTA de Blackwell.
  • La taille de grille doit être un multiple de la taille de cluster.
  • Cluster Launch Control (Blackwell) fournit une file de travail matérielle — la version câblée de ce que font les megakernels en logiciel.

Vérifiez que vous avez compris

Pourquoi les blocs d'un cluster doivent-ils être sur le même GPC ?

Parce que le réseau d'interconnexion SM-à-SM qui implémente DSMEM existe à l'intérieur d'un GPC. Deux SM de GPC différents n'ont pas de chemin direct : leur seul point de rendez-vous est le L2.

C'est aussi ce qui limite la taille de cluster : un GPC Hopper contient ~16-18 SM, ce qui borne naturellement le nombre de blocs qu'on peut y placer simultanément.

Un cluster de 4 blocs pour une GEMM. Le trafic HBM total est-il divisé par 4 ?

Non. Le multicast ne s'applique qu'aux opérandes partagés entre les blocs du cluster.

Dans un cluster de 4 blocs disposés en ligne (même bande de lignes de \(\mathbf{C}\)), la tuile de \(\mathbf{A}\) est partagée par les 4 : trafic de \(\mathbf{A}\) divisé par 4. Mais chaque bloc a sa propre tuile de \(\mathbf{B}\) : trafic de \(\mathbf{B}\) inchangé.

Si \(\mathbf{A}\) et \(\mathbf{B}\) contribuent également, le gain global est de \(\left(1 - \frac{1}{2}\left(1 - \frac{1}{4}\right)\right) = 62{,}5\ \%\) du trafic initial, soit 37,5 % d'économie.

Avec une disposition 2×2, on partage \(\mathbf{A}\) entre 2 et \(\mathbf{B}\) entre 2, soit 50 % d'économie sur les deux — souvent le meilleur choix.

En quoi Cluster Launch Control ressemble-t-il au runtime d'un megakernel ?

Les deux résolvent le même problème : affecter dynamiquement du travail à des unités persistantes.

Dans Mirage MPK, des SM « ordonnanceurs » distribuent les tâches aux SM « workers » via des files, avec un mélange de distribution juste-à-temps (pour l'équilibrage) et anticipée (pour la latence). Chez Hazy Research, le megakernel de débit utilise une file de travail globale qui apporte 14,2 % d'amélioration sur un ordonnancement en tourniquet à lot 8 192.

CLC fait la même chose en matériel, sans coût d'atomiques ni de SM dédiés. C'est un signe que les concepteurs de puces suivent l'évolution des logiciels — et une raison de penser que les megakernels vont devenir plus faciles à écrire, pas moins.

Chapitre suivant : 4 · wgmma et tcgen05

Sources de ce chapitre