Aller au contenu

2 · Coalescence et conflits de banc

Deux façons de gaspiller une ressource sans que rien ne le signale. Le résultat est juste, le programme tourne, et vous perdez un facteur 8 à 32.

2.1 La coalescence, mécaniquement

Quand les 32 threads d'un warp exécutent une instruction de chargement, le matériel :

  1. collecte les 32 adresses demandées ;
  2. les regroupe en secteurs de 32 octets ;
  3. émet une transaction mémoire par secteur distinct.

L'efficacité est le rapport entre octets utiles et octets transférés :

\[ \eta = \frac{\text{octets demandés par le warp}}{32 \times \text{nombre de secteurs}} \]

Le cas parfait

int i = blockIdx.x * blockDim.x + threadIdx.x;
float x = donnees[i];

32 threads × 4 octets = 128 octets contigus = 4 secteurs.

\[\eta = \frac{128}{4 \times 32} = 100\ \%\]

Le cas à pas 2

float x = donnees[i * 2];

Les adresses vont de 0 à 248 avec un pas de 8 octets : 256 octets d'étendue = 8 secteurs pour 128 octets utiles.

\[\eta = \frac{128}{8 \times 32} = 50\ \%\]

Le cas dispersé

float x = donnees[indices[i]];   // indices aléatoires

Chaque thread tombe dans un secteur différent : 32 secteurs pour 128 octets.

\[\eta = \frac{128}{32 \times 32} = 12{,}5\ \%\]

Le tableau à retenir

Motif d'accès Secteurs par warp Efficacité
Contigu aligné (a[i]) 4 100 %
Contigu non aligné 5 80 %
Pas de 2 (a[2i]) 8 50 %
Pas de 4 (a[4i]) 16 25 %
Pas ≥ 8 32 12,5 %
Aléatoire ~32 12,5 %
Diffusion (a[0] partout) 1 100 %*

* La diffusion est un cas particulier optimisé par le matériel : une seule transaction, valeur répliquée.

2.2 Les cinq causes de non-coalescence, et leurs remèdes

Cause 1 — l'ordre des dimensions inversé

Le classique absolu.

// MAUVAIS : x indexe les lignes
int lig = blockIdx.x * blockDim.x + threadIdx.x;
int col = blockIdx.y * blockDim.y + threadIdx.y;
float v = matrice[lig * N + col];   // threads voisins → lignes voisines
                                    // → adresses espacées de N*4 octets

// BON : x indexe les colonnes (dimension contiguë en row-major)
int col = blockIdx.x * blockDim.x + threadIdx.x;
int lig = blockIdx.y * blockDim.y + threadIdx.y;
float v = matrice[lig * N + col];   // threads voisins → colonnes voisines

Symptôme au profileur : L1/TEX Global Load Access Pattern signale un nombre excessif de secteurs par requête (8× ou 32× au lieu de 4).

Cause 2 — la disposition AoS

Array of Structures contre Structure of Arrays :

// AoS : chaque particule est un bloc de 24 octets
struct Particule { float x, y, z, vx, vy, vz; };
Particule* p;
float px = p[i].x;    // pas de 24 octets → η ≈ 17 %

// SoA : chaque champ est un tableau contigu
struct Particules { float *x, *y, *z, *vx, *vy, *vz; };
float px = p.x[i];    // contigu → η = 100 %

La règle SoA

Sur GPU, toujours préférer Structure of Arrays. C'est l'inverse de l'intuition CPU (où l'AoS optimise la localité pour un accès à toute la structure), et c'est une des adaptations mentales les plus importantes du portage CPU → GPU.

Cas intermédiaire utile : AoSoA (array of structures of arrays), où l'on groupe par blocs de 32 : struct { float x[32], y[32], z[32]; }. On obtient la coalescence et la localité de structure.

Cause 3 — le désalignement

float* base = donnees + 1;      // décalage de 4 octets
float x = base[i];              // → 5 secteurs au lieu de 4, η = 80 %

cudaMalloc renvoie toujours une adresse alignée sur 256 octets. Les désalignements viennent des sous-vues et des décalages calculés. Ils coûtent 20 %, ce qui est modéré mais gratuit à éviter.

Pour les accès vectorisés (float4), le désalignement n'est pas une perte de performance mais une erreur :

float4 v = *reinterpret_cast<float4*>(ptr);   // exige ptr % 16 == 0

Cause 4 — l'indirection

float x = valeurs[indices[i]];

Inévitable dans les graphes, les MoE, les tables de plongement. Trois atténuations :

  • trier les indices avant, pour que les threads voisins accèdent à des zones voisines (cub::DeviceRadixSort) ;
  • regrouper par bloc : trier les jetons par expert avant le GEMM du MoE, c'est exactement ce que fait la fusion gather-GEMM de Mirage MPK, qui élimine un prétraitement représentant « jusqu'à 11 % du temps d'exécution du MoE » ;
  • exploiter le cache : si la table est petite (< 50 Mo), elle tient en L2 et l'indirection devient acceptable.

Cause 5 — la transposition

Lire une matrice par colonnes est intrinsèquement non coalescé. Le remède est le passage par la mémoire partagée :

__global__ void transposer(const float* in, float* out, int N) {
    __shared__ float tuile[32][33];   // 33 pour éviter les conflits de banc

    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;

    // Lecture coalescée depuis in
    if (x < N && y < N) tuile[threadIdx.y][threadIdx.x] = in[y * N + x];
    __syncthreads();

    // Écriture coalescée dans out (indices de bloc échangés)
    x = blockIdx.y * 32 + threadIdx.x;
    y = blockIdx.x * 32 + threadIdx.y;
    if (x < N && y < N) out[y * N + x] = tuile[threadIdx.x][threadIdx.y];
}

Les deux accès à la mémoire globale sont coalescés. La transposition a lieu en mémoire partagée, où les accès dispersés coûtent beaucoup moins cher — à condition d'éviter les conflits de banc, d'où le [33].

2.3 Les conflits de banc

La mémoire partagée est découpée en 32 bancs de 4 octets, entrelacés :

\[ \text{banc}(a) = \left\lfloor \frac{a}{4} \right\rfloor \bmod 32 \]

où \(a\) est l'adresse en octets.

Un warp accède à la mémoire partagée en un cycle si :

  • les 32 threads touchent 32 bancs distincts, ou
  • plusieurs threads touchent la même adresse (diffusion, gérée par le matériel).

Sinon, l'accès est sérialisé : un conflit à \(k\) voies coûte \(k\) cycles.

Le calcul du degré de conflit

Pour un accès partage[f(tid)] en float, le banc du thread \(t\) est \(f(t) \bmod 32\). Le degré de conflit est le nombre maximal de threads partageant un même banc avec des adresses différentes.

Accès Banc du thread \(t\) Conflit
s[t] \(t\) aucun
s[2t] \(2t \bmod 32\) 2 voies
s[4t] \(4t \bmod 32\) 4 voies
s[32t] \(0\) 32 voies
s[33t] \(33t \bmod 32 = t\) aucun
s[0] \(0\) aucun (diffusion)

La ligne s[33t] explique le padding : ajouter 1 à la largeur d'un tableau 2D change le pas de 32 à 33, et \(33 \equiv 1 \pmod{32}\).

Le cas des types larges

Un double (8 octets) occupe deux bancs. Un accès s[t] en double fait donc que les threads 0 et 16 touchent le même banc. Le matériel gère ce cas spécialement (il découpe la transaction en deux phases), mais il faut le savoir pour interpréter le profileur.

Un float4 (16 octets) est traité en quatre phases. C'est encore efficace, mais le calcul des conflits doit se faire par phase.

2.4 Le swizzling

Le padding a deux défauts : il gaspille de la mémoire partagée (ressource rare), et il casse l'alignement requis par TMA et wgmma, qui exigent des tuiles contiguës de taille précise.

La solution moderne est le swizzling : une permutation XOR des adresses.

\[ a' = a \oplus \big( (a \gg s) \ \&\ m \big) \]

Pour un tableau \(32 \times 32\) de float, avec le swizzle classique \(a' = a \oplus ((a \gg 5) \,\&\, 31)\), la ligne \(r\) colonne \(c\) est stockée à \(32r + (c \oplus r)\). Un accès en colonne (c fixe, r variable) donne les bancs \((c \oplus r) \bmod 32\), tous distincts quand \(r\) varie.

Aucun octet gaspillé, aucun conflit, alignement préservé.

Qui fait ça

Vous n'écrivez normalement pas de swizzle à la main. C'est encodé dans les types de layout de CuTe (Swizzle<B,M,S>), dans ThunderKittens, et généré par Triton. Le comprendre est utile pour lire ces bibliothèques et pour déboguer.

2.5 Comment les détecter au profileur

Non-coalescence

Dans Nsight Compute, section Memory Workload Analysis :

Métrique Interprétation
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum ÷ requêtes secteurs par requête : 4 est parfait, 32 est catastrophique
Global Load Efficiency (héritée de nvprof) directement \(\eta\) en %
Excessive sectors dans les règles Nsight le signale explicitement

Nsight Compute émet aussi des règles textuelles :

[Warning] The memory access pattern for global loads might not be optimal.
On average, only 8.0 of the 32 bytes transmitted per sector are utilized.

Conflits de banc

Métrique Interprétation
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum nombre de conflits en lecture
..._op_st.sum idem en écriture
Shared Memory Bank Conflicts dans le rapport résumé lisible

Une valeur non nulle mais faible peut être acceptable. Une valeur du même ordre que le nombre d'accès signifie une sérialisation systématique.

2.6 Une méthode de diagnostic en trois étapes

Étape 1 — lire le code. Pour chaque accès mémoire dans une boucle chaude, posez-vous : « quand threadIdx.x augmente de 1, de combien augmente l'adresse ? »

  • +4 octets (ou +taille de l'élément) → coalescé ;
  • +0 → diffusion, bien aussi ;
  • autre chose → problème.

Étape 2 — calculer le trafic théorique. \(Q_{\text{utile}}\) contre \(Q_{\text{réel}} = Q_{\text{utile}} / \eta\). Si \(\eta = 12{,}5\ \%\), votre noyau transfère 8× plus que nécessaire.

Étape 3 — mesurer. dram__bytes.sum dans Nsight Compute donne les octets réellement lus en HBM. Comparez à votre \(Q_{\text{utile}}\). Le rapport est \(1/\eta\).

Résumé du chapitre

À retenir

  • Un warp émet une transaction par secteur de 32 octets touché. L'efficacité \(\eta\) peut descendre à 12,5 %.
  • Les cinq causes : ordre des dimensions inversé, disposition AoS, désalignement, indirection, transposition.
  • Sur GPU : Structure of Arrays, toujours.
  • La mémoire partagée a 32 bancs de 4 octets. Le padding ([33] au lieu de [32]) résout les conflits pour 128 octets gaspillés.
  • Le swizzling fait la même chose gratuitement et préserve l'alignement exigé par TMA et wgmma. C'est ce qu'utilisent CuTe et ThunderKittens.
  • Diagnostic : comparez dram__bytes.sum à votre volume utile théorique. Le rapport est \(1/\eta\).

Vérifiez que vous avez compris

Un noyau lit une matrice colonne par colonne. Vous mesurez 420 Go/s sur un H100. Est-ce cohérent avec un accès non coalescé ?

Oui, exactement. \(3\,350 \times 0{,}125 = 419\) Go/s.

Un accès par colonne dans une matrice row-major a un pas égal à la largeur de la ligne. Dès que ce pas dépasse 8 float (32 octets), chaque thread tombe dans un secteur différent : \(\eta = 12{,}5\ \%\).

Le remède est la transposition par tuiles en mémoire partagée, ou le changement de disposition en amont si vous contrôlez la production des données.

Pourquoi __shared__ double tab[32][32] avec un accès tab[tid][0] pose-t-il un conflit même après padding à [33] ?

Un double fait 8 octets et occupe deux bancs consécutifs. L'adresse en octets de tab[t][0] est \(8 \times 33 t = 264t\), donc le banc est \((264t/4) \bmod 32 = 66t \bmod 32 = 2t \bmod 32\).

Pour \(t\) et \(t+16\) : \(2t\) et \(2t + 32 \equiv 2t\). Conflit à 2 voies.

Pour un tableau de double, le bon padding est un nombre impair de bancs, soit ici [32][33] en double donne un pas de 66 bancs — pair. Il faudrait un pas donnant un nombre impair de bancs, par exemple en ajoutant 2 éléments : pas de \(34 \times 2 = 68\) bancs, toujours pair.

En pratique, pour les types 8 octets, on utilise le swizzling ou on restructure. C'est un bon exemple de la raison pour laquelle le padding naïf n'est pas une solution générale.

Vous avez une table de plongement de 128 000 lignes × 4 096 dimensions en BF16 (1 Go). Les accès sont indirects. Que faire ?

Le volume dépasse largement le L2 (50 Mo), donc l'indirection coûte plein tarif au niveau des lignes. Mais l'important est que chaque ligne accédée est lue en entier et est contiguë : 4 096 × 2 = 8 192 octets, soit 256 secteurs pleinement utilisés.

L'indirection porte donc sur le choix de la ligne, pas sur la lecture à l'intérieur. Assignez un bloc (ou un warp) par ligne et lisez-la de manière coalescée. L'efficacité de coalescence est alors de 100 %, et le seul coût de l'indirection est l'absence de localité entre lignes — c'est-à-dire l'impossibilité de préchargement, pas un gaspillage de bande passante.

C'est un point souvent mal compris : l'indirection sur des blocs larges n'est pas un problème de coalescence.

Chapitre suivant : 3 · Occupancy et latence

Sources de ce chapitre