Aller au contenu

1 · Asynchronisme et mbarrier

Comment charger la tuile \(n+1\) pendant qu'on calcule la tuile \(n\). C'est le pipelining logiciel, et c'est la technique qui sépare un noyau à 40 % du pic d'un noyau à 90 %.

1.1 Le problème du chargement synchrone

Reprenons la GEMM par tuiles de la partie 2 :

for (int t = 0; t < N/T; ++t) {
    As[ty][tx] = A[...];      // ← chargement bloquant : 400-800 cycles
    Bs[ty][tx] = B[...];
    __syncthreads();
    for (int k = 0; k < T; ++k) acc += As[ty][k] * Bs[k][tx];
    __syncthreads();
}

La chronologie d'un warp :

[charger 500 cy][sync][calculer 200 cy][sync][charger 500 cy][sync][calculer]…

Plus de 70 % du temps est passé à attendre. Avec assez de warps résidents, le SM reste occupé — mais on paie cette occupation en registres, donc en taille de tuile, donc en intensité arithmétique.

L'objectif du pipelining : superposer les deux phases.

warp : [charger t=0][charger t=1 ∥ calculer t=0][charger t=2 ∥ calculer t=1]…

1.2 Avant Ampere : le double tampon manuel

La technique classique, qui fonctionne partout :

__shared__ float As[2][T][T];    // deux tampons
__shared__ float Bs[2][T][T];

int etage = 0;
// Préchargement du premier
As[0][ty][tx] = A[...]; Bs[0][ty][tx] = B[...];
__syncthreads();

for (int t = 0; t < N/T; ++t) {
    int suivant = etage ^ 1;

    // Charger le suivant DANS DES REGISTRES (non bloquant tant qu'on
    // ne lit pas la valeur)
    float a_reg = (t + 1 < N/T) ? A[...] : 0.0f;
    float b_reg = (t + 1 < N/T) ? B[...] : 0.0f;

    // Calculer sur l'étage courant
    for (int k = 0; k < T; ++k) acc += As[etage][ty][k] * Bs[etage][k][tx];

    // Publier le suivant
    As[suivant][ty][tx] = a_reg;
    Bs[suivant][ty][tx] = b_reg;
    __syncthreads();

    etage = suivant;
}

L'astuce : a_reg = A[...] émet la requête mémoire mais ne bloque que lorsqu'on utilise a_reg, c'est-à-dire après la boucle de calcul. Le matériel a un scoreboard qui suit les dépendances.

Deux défauts :

  1. la donnée transite par les registres — elle occupe des registres pendant toute la latence, ce qui limite la taille de tuile ;
  2. on ne peut avoir que deux étages en pratique, faute de registres.

1.3 Ampere : cp.async

Ampere introduit une instruction qui copie directement de la mémoire globale vers la mémoire partagée, sans passer par les registres.

En PTX :

cp.async.cg.shared.global [%r_smem], [%r_gmem], 16;
cp.async.commit_group;
cp.async.wait_group 1;      // attendre tous les groupes sauf le dernier

En CUDA C++, via l'API cuda::pipeline :

#include <cuda/pipeline>
#include <cooperative_groups.h>
namespace cg = cooperative_groups;

template <int ETAGES>
__global__ void gemm_pipeline(const float* A, const float* B, float* C, int N) {
    __shared__ float As[ETAGES][T][T];
    __shared__ float Bs[ETAGES][T][T];

    auto bloc = cg::this_thread_block();
    __shared__ cuda::pipeline_shared_state<cuda::thread_scope_block, ETAGES> etat;
    auto pipe = cuda::make_pipeline(bloc, &etat);

    // Amorçage : remplir ETAGES-1 étages
    for (int t = 0; t < ETAGES - 1; ++t) {
        pipe.producer_acquire();
        cuda::memcpy_async(bloc, &As[t][0][0], &A[/* ... */],
                           sizeof(float) * T * T, pipe);
        cuda::memcpy_async(bloc, &Bs[t][0][0], &B[/* ... */],
                           sizeof(float) * T * T, pipe);
        pipe.producer_commit();
    }

    // Boucle principale
    for (int t = 0; t < N / T; ++t) {
        int etage = t % ETAGES;

        pipe.consumer_wait();                     // attendre l'étage courant
        for (int k = 0; k < T; ++k)
            acc += As[etage][ty][k] * Bs[etage][k][tx];
        pipe.consumer_release();

        // Émettre le chargement de t + ETAGES - 1
        if (t + ETAGES - 1 < N / T) {
            pipe.producer_acquire();
            cuda::memcpy_async(bloc, &As[etage][0][0], &A[/* ... */],
                               sizeof(float) * T * T, pipe);
            pipe.producer_commit();
        }
    }
}

Avantages sur le double tampon manuel :

Registres Double tampon cp.async
Chemin global → reg → shared idem global → shared
Registres consommés élevé élevé nul
Étages possibles 2 2 3 à 6
Bypass du L1 non non oui (variante .cg)

Combien d'étages ?

Le nombre d'étages doit couvrir la latence mémoire :

\[\text{étages} \ge \frac{\text{latence mémoire}}{\text{temps de calcul par étage}} + 1\]

Avec 500 cycles de latence et 200 cycles de calcul par tuile : 3,5, donc 4 étages. Le coût est en mémoire partagée : 4 × 2 × 32 × 32 × 4 = 32 Ko.

C'est le compromis central du pipelining : plus d'étages cachent mieux la latence mais consomment la mémoire partagée qui limite l'occupancy.

1.4 Hopper : mbarrier

cp.async avec wait_group a une limite : la granularité est le groupe, et l'attente est ordonnée (on attend « les \(n\) derniers groupes »). Pour un pipeline sophistiqué — surtout avec des producteurs et consommateurs distincts — il faut mieux.

La mbarrier (memory barrier) de Hopper est un objet de synchronisation de 8 octets résidant en mémoire partagée, avec un compteur de participants et un compteur d'octets attendus.

// Initialisation : nb_threads participants
mbarrier.init.shared.b64 [barriere], 128;

// Le producteur annonce combien d'octets vont arriver
mbarrier.arrive.expect_tx.shared.b64 _, [barriere], 16384;

// … la copie TMA écrit et incrémente automatiquement le compteur d'octets …

// Le consommateur attend
mbarrier.try_wait.parity.shared.b64 %p, [barriere], %phase;

Trois propriétés qui la distinguent de __syncthreads() :

  1. elle compte des octets, pas seulement des arrivées. Une copie TMA incrémente le compteur du nombre d'octets transférés, ce qui permet d'attendre « la donnée est arrivée » et non « le thread a fini d'émettre » ;
  2. elle a une phase (un bit qui bascule à chaque cycle complet), ce qui permet de la réutiliser dans une boucle sans réinitialisation ;
  3. elle est asynchrone : try_wait peut échouer et rendre la main, ce qui permet d'attendre en faisant autre chose.

En CUDA C++ :

#include <cuda/barrier>

__global__ void exemple() {
    __shared__ cuda::barrier<cuda::thread_scope_block> barriere;

    if (threadIdx.x == 0) {
        init(&barriere, blockDim.x);
    }
    __syncthreads();

    // ... travail ...

    barriere.arrive_and_wait();
}

1.5 Le motif complet : pipeline à N étages avec mbarrier

C'est la structure de tout noyau Hopper moderne. Voici le squelette conceptuel.

constexpr int ETAGES = 4;

__shared__ alignas(128) Tuile tampon[ETAGES];
__shared__ cuda::barrier<cuda::thread_scope_block> plein[ETAGES];
__shared__ cuda::barrier<cuda::thread_scope_block> vide[ETAGES];

// Initialisation par un seul thread
if (threadIdx.x == 0) {
    for (int i = 0; i < ETAGES; ++i) {
        init(&plein[i], 1);            // 1 producteur
        init(&vide[i],  NB_CONSOMMATEURS);
    }
}
__syncthreads();

// ── Warp producteur ────────────────────────────────
if (est_producteur) {
    for (int t = 0; t < nb_tuiles; ++t) {
        int e = t % ETAGES;
        vide[e].wait(vide[e].arrive());       // attendre que l'étage soit libre
        lancer_copie(&tampon[e], source(t), plein[e]);
    }
}

// ── Warps consommateurs ────────────────────────────
else {
    for (int t = 0; t < nb_tuiles; ++t) {
        int e = t % ETAGES;
        plein[e].wait(/* phase */);            // attendre les données
        calculer(tampon[e]);
        vide[e].arrive();                      // libérer l'étage
    }
}

Deux jeux de barrières : plein[e] signale « les données de l'étage \(e\) sont arrivées », vide[e] signale « l'étage \(e\) peut être réécrit ». C'est un tampon circulaire classique, transposé sur GPU.

Pourquoi c'est la bonne structure

Elle sépare complètement le mouvement des données du calcul. Le producteur n'attend jamais le calcul, le consommateur n'attend jamais la mémoire au-delà de son étage.

C'est aussi la structure exacte des megakernels : dans MPK, les « workers » sont les consommateurs, les « schedulers » gèrent les dépendances, et les tâches sont les tuiles. Chez Hazy Research, les instructions sont pipelinées de la même façon, avec un allocateur de pages qui joue le rôle des étages.

1.6 Le coût des barrières

Un point mesuré et rarement mentionné : une barrière asynchrone coûte ~60 nanosecondes.

Dans le décompte de Hazy Research pour une passe avant de Llama-1B sur B200 (600 µs au total), 40 µs sont attribués à la « synchronisation de bas niveau entre warps », soit environ 6,7 % du temps total.

Ce n'est pas énorme, mais c'est le genre de coût qui devient dominant quand tout le reste a été optimisé. Conséquence pratique : ne synchronisez pas plus finement que nécessaire. Un pipeline à 8 étages avec des barrières par tuile de 4 Ko passe plus de temps à se synchroniser qu'à transférer.

1.7 Récapitulatif des mécanismes

Mécanisme Génération Ce qu'il apporte Limite
Double tampon en registres toutes pipelining basique consomme des registres, 2 étages
cp.async Ampere+ global → shared direct attente par groupes ordonnés
cuda::pipeline Ampere+ API C++ sur cp.async idem
mbarrier Hopper+ attente par octets, phases pénible à écrire
cuda::barrier Hopper+ API C++ sur mbarrier moins expressive que le PTX
TMA Hopper+ copies multi-D par un seul thread chapitre suivant

Résumé du chapitre

À retenir

  • Le pipelining superpose le chargement de la tuile \(n+1\) au calcul de la tuile \(n\). C'est ce qui permet de saturer la mémoire avec peu de warps.
  • cp.async (Ampere) copie global → partagé sans passer par les registres, ce qui permet 3 à 6 étages au lieu de 2.
  • Nombre d'étages ≈ latence / temps de calcul par étage + 1. Le coût est en mémoire partagée.
  • mbarrier (Hopper) compte des octets et a une phase : c'est ce qui permet d'attendre l'arrivée effective des données d'une copie asynchrone.
  • Le motif canonique est un tampon circulaire avec deux jeux de barrières (plein / vide), producteur et consommateurs séparés.
  • Une barrière coûte ~60 ns ; ne synchronisez pas plus finement que nécessaire.

Vérifiez que vous avez compris

Pourquoi cp.async permet-il plus d'étages de pipeline que le double tampon en registres ?

Parce que la donnée en vol n'occupe aucun registre. Avec le double tampon manuel, chaque valeur en transit occupe un registre pendant toute la latence mémoire (500 cycles) ; avec une tuile de 32×32 float et 256 threads, cela fait 4 registres par thread par étage — vite prohibitif.

Avec cp.async, seule la mémoire partagée est consommée, et elle est beaucoup plus abondante (227 Ko contre 256 Ko de registres pour 2 048 threads, soit 128 octets par thread).

Pourquoi une mbarrier a-t-elle besoin d'un compteur d'octets en plus d'un compteur d'arrivées ?

Parce qu'avec TMA, un seul thread émet la copie mais la donnée arrive plus tard, de façon asynchrone. Un compteur d'arrivées classique compterait « 1 thread est passé », ce qui ne dit rien sur la présence des données.

Le compteur d'octets (transaction count) est incrémenté par le matériel au fur et à mesure de l'écriture en mémoire partagée. Le consommateur attend donc « les 16 384 octets attendus sont arrivés », ce qui est la vraie condition.

D'où le expect_tx : le producteur annonce combien d'octets vont arriver avant de lancer la copie.

Votre pipeline à 6 étages est plus lent qu'à 4 étages. Explication plausible ?

Deux causes, toutes deux fréquentes :

  1. La mémoire partagée. 6 étages consomment 1,5× plus que 4. Si cela fait passer le nombre de blocs résidents de 2 à 1, l'occupancy est divisée par deux et le gain du pipelining est effacé.
  2. Le coût des barrières. À 60 ns par barrière et 2 barrières par étage par tuile, un pipeline profond avec de petites tuiles passe une fraction croissante de son temps à se synchroniser.

La bonne démarche : balayer {2, 3, 4, 5, 6} et mesurer. L'optimum dépend de la taille de tuile, de la carte et du noyau.

Chapitre suivant : 2 · Le Tensor Memory Accelerator

Sources de ce chapitre