Aller au contenu

7 · Atomiques et synchronisation

Le chapitre qui ferme la partie 2 et ouvre directement la partie 8. Comment des threads qui ne se connaissent pas peuvent coopérer — et à quel prix.

7.1 Les opérations atomiques

Une opération atomique lit, modifie et écrit une adresse sans que personne puisse s'intercaler.

int  ancien = atomicAdd(&compteur, 1);       // renvoie l'ancienne valeur
float f     = atomicAdd(&somme, 3.14f);      // flottant : Pascal+
int  m      = atomicMax(&maximum, valeur);
int  e      = atomicExch(&cellule, nouvelle);
int  c      = atomicCAS(&verrou, attendu, nouveau);   // compare-and-swap

atomicCAS est le primitif universel : tous les autres peuvent s'en déduire.

// atomicMax pour float, qui n'existe pas nativement
__device__ float atomicMaxFloat(float* adr, float val) {
    int* adr_i = (int*)adr;
    int ancien = *adr_i, suppose;
    do {
        suppose = ancien;
        ancien  = atomicCAS(adr_i, suppose,
                            __float_as_int(fmaxf(val, __int_as_float(suppose))));
    } while (suppose != ancien);
    return __int_as_float(ancien);
}

Ce truc ne marche que pour les positifs

La comparaison d'entiers ne correspond à celle des flottants que pour les valeurs positives (la représentation IEEE 754 des flottants positifs est monotone quand on la lit comme un entier). Pour des valeurs signées, il faut un traitement du bit de signe. C'est un piège très classique.

Où vivent les atomiques

Portée Où Latence
Mémoire partagée dans le SM ~30 cycles
Mémoire globale au niveau du L2 ~200-400 cycles
Système (atomicAdd_system) cohérent avec le CPU et les autres GPU très lent

Point important : un atomique en mémoire globale s'exécute dans le L2, pas dans le SM. C'est ce qui le rend visible par tous les blocs, et c'est aussi ce qui le rend lent.

7.2 Le coût réel : la contention

Le débit d'un atomique dépend entièrement du nombre de threads qui visent la même adresse.

32 threads → 32 adresses différentes   : ~1 transaction, rapide
32 threads → 1 adresse identique       : sérialisé, 32× plus lent

Trois stratégies pour réduire la contention, par ordre d'efficacité.

1. Agréger dans le warp (vu au chapitre 4) :

// Au lieu de 32 atomicAdd sur la même adresse
float somme_warp = reduce_add_warp(ma_valeur);
if ((threadIdx.x & 31) == 0) atomicAdd(&total, somme_warp);

2. Privatiser en mémoire partagée. Le motif de l'histogramme :

__global__ void histogramme(const unsigned char* donnees, int n,
                            unsigned int* hist_global) {
    __shared__ unsigned int hist_local[256];

    // 1. Initialiser la copie privée
    for (int i = threadIdx.x; i < 256; i += blockDim.x) hist_local[i] = 0;
    __syncthreads();

    // 2. Accumuler dans la copie privée (atomiques en mémoire partagée)
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n; i += blockDim.x * gridDim.x) {
        atomicAdd(&hist_local[donnees[i]], 1);
    }
    __syncthreads();

    // 3. Fusionner : 256 atomiques globaux par bloc, pas n
    for (int i = threadIdx.x; i < 256; i += blockDim.x) {
        if (hist_local[i]) atomicAdd(&hist_global[i], hist_local[i]);
    }
}

Gain typique : facteur 10 à 50 selon la distribution des données.

3. Répliquer les compteurs. Si même la mémoire partagée sature (données très concentrées sur peu de valeurs), on maintient \(k\) copies de l'histogramme et chaque warp écrit dans la copie warp_id % k.

7.3 Le modèle mémoire de CUDA

Voici la partie que presque personne ne lit et qui décide de la correction des megakernels.

Le problème

// Bloc producteur
donnees[0] = 42;
pret = 1;

// Bloc consommateur
while (pret == 0) { }
int x = donnees[0];    // ← garanti d'être 42 ?

Non, pas sans annotation. Trois choses peuvent casser :

  1. le compilateur peut réordonner les deux écritures, ou garder pret dans un registre et boucler à l'infini ;
  2. le matériel peut rendre pret = 1 visible avant donnees[0] = 42 ;
  3. les caches ne sont pas cohérents entre L1 de SM différents.

La solution : le modèle mémoire C++ appliqué à CUDA

Depuis CUDA 9, CUDA implémente un modèle mémoire formel inspiré de celui de C++11, exposé par cuda::atomic :

#include <cuda/atomic>

// Producteur
donnees[0] = 42;
pret.store(1, cuda::memory_order_release);

// Consommateur
while (pret.load(cuda::memory_order_acquire) == 0) { }
int x = donnees[0];    // ← maintenant garanti

La paire release/acquire garantit que tout ce qui a été écrit avant le release est visible après le acquire correspondant.

Il faut aussi préciser la portée :

cuda::atomic<int, cuda::thread_scope_block>   compteur_bloc;
cuda::atomic<int, cuda::thread_scope_device>  compteur_gpu;    // ← inter-blocs
cuda::atomic<int, cuda::thread_scope_system>  compteur_systeme;

Plus la portée est large, plus l'opération est coûteuse. Utiliser thread_scope_device là où thread_scope_block suffirait est un gaspillage courant.

La version bas niveau

En CUDA C++ classique, on utilise volatile + __threadfence() :

volatile int* pret_v = &pret;
donnees[0] = 42;
__threadfence();          // toutes les écritures précédentes sont visibles
*pret_v = 1;

volatile empêche le compilateur de mettre la variable en cache dans un registre. __threadfence() est la barrière mémoire au niveau du périphérique.

volatile seul ne suffit jamais

volatile traite le problème du compilateur, pas celui du matériel ni celui des caches. Beaucoup de code trouvé en ligne utilise volatile sans __threadfence() : il fonctionne souvent par chance, et casse à la première optimisation matérielle.

L'écriture correcte en 2026 utilise cuda::atomic avec un ordre mémoire explicite.

7.4 La barrière inter-blocs, écrite à la main

Ce code est le lien direct avec les megakernels. Il est correct uniquement si tous les blocs sont simultanément résidents.

#include <cuda/atomic>

__device__ cuda::atomic<unsigned, cuda::thread_scope_device> arrives{0};
__device__ cuda::atomic<unsigned, cuda::thread_scope_device> generation{0};

__device__ void barriere_grille(unsigned nb_blocs) {
    __syncthreads();                     // d'abord synchroniser le bloc

    if (threadIdx.x == 0) {
        unsigned gen = generation.load(cuda::memory_order_relaxed);

        if (arrives.fetch_add(1, cuda::memory_order_acq_rel) == nb_blocs - 1) {
            // dernier arrivé : réinitialiser et libérer tout le monde
            arrives.store(0, cuda::memory_order_relaxed);
            generation.fetch_add(1, cuda::memory_order_release);
        } else {
            // attendre le changement de génération
            while (generation.load(cuda::memory_order_acquire) == gen) { }
        }
    }
    __syncthreads();                     // rediffuser au bloc
}

La condition de résidence n'est pas négociable

Si vous lancez plus de blocs que le GPU ne peut en héberger simultanément, ce code interblo­que : les blocs résidents attendent des blocs qui ne seront jamais lancés, parce que l'ordonnanceur attend qu'un SM se libère.

C'est pourquoi cooperative_groups::grid_group::sync() exige cudaLaunchCooperativeKernel et une grille dimensionnée par cudaOccupancyMaxActiveBlocksPerMultiprocessor. C'est aussi pourquoi tout megakernel est un noyau persistant : exactement une vague de blocs, qui ne se termine jamais avant la fin du travail.

Le compteur de génération évite le bug classique où un bloc rapide franchit la barrière, revient et incrémente arrives avant que les lents aient vu 0.

7.5 Du verrou global aux dépendances fines

La barrière de la section précédente coûte cher. Rappel des mesures de Kog sur MI300X :

Approche Latence de synchronisation
Barrière atomique naïve 7,59 – 7,88 µs
Sentinelle NaN 0,80 – 0,93 µs

La technique de la sentinelle est élégante : au lieu d'un compteur séparé, on initialise le tampon de sortie avec une valeur impossible (ici NaN). Le consommateur lit la donnée elle-même et attend qu'elle ne soit plus NaN. Une seule lecture au lieu de deux, et aucune opération atomique.

// Producteur : écrit simplement la donnée
sortie[i] = valeur;                     // remplace le NaN

// Consommateur : attend que la donnée soit valide
float v;
do { v = sortie[i]; } while (isnan(v));

Les limites de la sentinelle

  • Il faut une valeur impossible dans le domaine de sortie. NaN convient pour des activations, pas pour des données arbitraires.
  • Il faut réinitialiser le tampon à chaque itération, ce qui coûte de la bande passante.
  • La visibilité mémoire reste à garantir : la lecture doit contourner le L1 (ld.global.cv ou volatile), sinon le consommateur peut lire indéfiniment une valeur mise en cache.

Mais l'idée générale est le vrai enseignement : on ne synchronise pas tout le monde, on attend exactement la donnée dont on a besoin.

C'est le principe du megakernel. Là où une frontière de noyau dit « attends que tout soit fini », un megakernel dit « attends que le compteur 47 atteigne 3 ».

Hazy Research le formule ainsi : des compteurs en mémoire globale, initialisés à zéro avant le lancement, incrémentés à la fin de chaque instruction ; les instructions dépendantes attendent que leurs compteurs atteignent leur valeur cible. Et leur exemple canonique : plutôt que d'attendre tout l'état caché du MLP avant la projection descendante, on le produit et le consomme en quatre morceaux avec quatre compteurs.

Mirage MPK généralise cela avec une notion d'événement : chaque tâche a au plus un événement déclencheur et un événement de sortie, et le compilateur fusionne les événements redondants.

7.6 Récapitulatif des mécanismes de synchronisation

Mécanisme Portée Coût Quand
__syncwarp() warp ~1 cycle après un accès partagé intra-warp
__syncthreads() bloc ~10-50 cycles motif charger/calculer
__syncthreads_count/and/or bloc idem + réduction vote sur le bloc
cuda::barrier / mbarrier bloc, asynchrone variable pipeline avec TMA
cluster.sync() cluster (Hopper+) ~100 cycles coopération multi-SM
grid.sync() grille µs rare, exige la résidence
atomic + attente arbitraire dépend dépendances fines, megakernels
Fin de noyau grille 1,3 à 10 µs le cas par défaut

La règle d'or

Synchronisez à la portée la plus étroite qui suffit. Un __syncthreads() là où un __syncwarp() suffit coûte 10 à 50× plus cher. Un grid.sync() là où un compteur dédié suffirait coûte 1 000× plus cher.

Résumé du chapitre

À retenir

  • Les atomiques globaux s'exécutent dans le L2. Leur coût dépend de la contention, pas de leur nombre.
  • Trois remèdes : agréger dans le warp, privatiser en mémoire partagée, répliquer les compteurs.
  • Sans annotation mémoire (cuda::atomic avec release/acquire, ou volatile + __threadfence()), la communication inter-blocs est incorrecte, même quand elle semble marcher.
  • Une barrière de grille écrite à la main n'est correcte que si tous les blocs sont résidents. D'où : tout megakernel est un noyau persistant.
  • Le passage de « barrière globale » à « attendre exactement la donnée dont j'ai besoin » est l'idée centrale de la partie 8.

Vérifiez que vous avez compris

Pourquoi un histogramme à 256 cases est-il beaucoup plus rapide qu'un histogramme à 2 cases, à volume de données égal ?

Parce que la contention est inversement proportionnelle au nombre de cases. Avec 2 cases, tous les threads visent l'une des deux adresses : les atomiques sont massivement sérialisés. Avec 256 cases et des données uniformes, la contention est divisée par 128.

Pour le cas à 2 cases (ou pour des données très concentrées), la bonne approche n'est pas l'atomique du tout : c'est une réduction. On compte localement dans chaque warp avec __ballot_sync + __popc, puis on agrège.

Ce code fonctionne sur votre machine et échoue sur une autre. Pourquoi ? 

__device__ int pret = 0;
// bloc 0
resultat[0] = calculer();
pret = 1;
// bloc 1
while (pret == 0) { }
utiliser(resultat[0]);
Trois raisons possibles, toutes réelles :

  1. pret n'est pas volatile : le compilateur peut le charger une fois en registre et boucler à l'infini. Le comportement dépend du niveau d'optimisation.
  2. Pas de __threadfence() : rien ne garantit que resultat[0] soit visible avant pret.
  3. Interblocage possible : si le bloc 1 est ordonnancé avant le bloc 0 et que le GPU est trop petit pour les héberger tous les deux, le bloc 0 ne démarre jamais.

La version correcte utilise cuda::atomic<int, thread_scope_device> avec release/acquire et garantit la résidence.

Dans un megakernel, pourquoi remplacer une barrière de grille par des compteurs fait-il gagner autant ?

Deux raisons cumulatives.

Coût unitaire : une barrière de grille fait communiquer tous les blocs entre eux (mesuré à ~7,6 µs sur MI300X en version naïve). Un compteur dédié ne fait communiquer que les blocs concernés.

Parallélisme conservé : la barrière impose que toute l'étape \(n\) soit finie avant toute l'étape \(n+1\). Avec des compteurs fins, le bloc qui produit le premier quart de l'état caché débloque immédiatement son consommateur, pendant que les autres quarts sont encore en cours. La profondeur du chemin critique diminue, pas seulement le coût par synchronisation.

Partie suivante : Performance

Sources de ce chapitre