Aller au contenu

5 · Les tensor cores

Une unité qui fait une multiplication matricielle entière en une instruction. C'est la raison pour laquelle un GPU de 2026 est 30× plus rapide qu'un GPU de 2016 sur l'apprentissage profond, alors que sa fréquence n'a pas bougé.

5.1 L'idée

Un cœur CUDA ordinaire exécute une multiplication-accumulation scalaire :

\[d = a \times b + c\]

Un tensor core exécute une multiplication-accumulation matricielle :

\[\mathbf{D} = \mathbf{A} \times \mathbf{B} + \mathbf{C}\]

où \(\mathbf{A}\), \(\mathbf{B}\), \(\mathbf{C}\), \(\mathbf{D}\) sont de petites matrices, typiquement \(16 \times 16 \times 16\) pour les premières générations.

Le gain n'est pas magique, il est combinatoire. Une MMA \(16\times16\times16\) effectue :

\[16 \times 16 \times 16 \times 2 = 8\,192 \text{ opérations flottantes}\]

(le facteur 2 compte une multiplication et une addition) tout en ne lisant que \(3 \times 16 \times 16 = 768\) éléments. L'intensité arithmétique de l'instruction elle-même est donc de \(8\,192 / 768 \approx 10{,}7\) opérations par élément lu, contre 1 pour une FMA scalaire.

Intuition

Un tensor core, c'est un réseau systolique miniature câblé dans le SM : une grille d'additionneurs-multiplieurs où les données circulent et se réutilisent sans repasser par les registres. Le silicium exploite la réutilisation intrinsèque du produit matriciel, que le logiciel devrait sinon organiser à la main.

5.2 Ce que ça coûte : l'accord sur les formats

Un tensor core n'accepte pas n'importe quoi. Chaque génération définit strictement :

  • les types d'entrée (FP16, BF16, TF32, FP8, FP4, INT8…) ;
  • le type d'accumulation (souvent FP32, parfois FP16 ou INT32) ;
  • les formes de tuiles supportées (\(m \times n \times k\)) ;
  • la disposition exacte des données dans les registres de chaque thread.

Ce dernier point est le plus pénible. Pour mma.sync.m16n8k16, chaque thread du warp doit détenir des fragments précis de \(\mathbf{A}\) et \(\mathbf{B}\) dans des registres nommés, selon une table de correspondance donnée dans le PTX ISA. Se tromper produit un résultat faux, pas une erreur.

C'est pour cette raison que presque personne n'écrit d'instruction MMA directement. On passe par :

  • wmma:: (l'API C++ de haut niveau, simple mais qui n'atteint pas le pic) ;
  • CUTLASS / CuTe, qui encode ces dispositions dans des types ;
  • Triton, qui les génère ;
  • ThunderKittens, qui les enveloppe dans des tuiles de 16×16.

5.3 L'évolution, génération par génération

Génération Instruction Portée Nouveauté
Volta (2017) hmma.884 8 threads (quadpair) l'invention
Turing (2018) mma.m16n8k8 warp INT8, INT4
Ampere (2020) mma.m16n8k16 warp BF16, TF32, sparsité 2:4, cp.async
Hopper (2022) wgmma warp group (128 threads) opérandes en mémoire partagée, asynchrone, FP8
Blackwell (2024) tcgen05.mma 1 ou 2 SM Tensor Memory, FP6/FP4, tuiles 256×256

Deux ruptures méritent qu'on s'y arrête.

Hopper : wgmma, l'opérande qui reste en mémoire partagée

Avant Hopper, les opérandes d'une MMA devaient être dans les registres des threads. Il fallait donc les charger explicitement depuis la mémoire partagée, avec l'instruction ldmatrix, en respectant la disposition exacte.

wgmma accepte que \(\mathbf{A}\) et \(\mathbf{B}\) soient directement en mémoire partagée, décrits par un descripteur de 64 bits. Le matériel va les chercher lui-même. Cela libère beaucoup de registres et permet des tuiles plus grandes.

wgmma est aussi asynchrone : on la lance, on fait autre chose, on attend avec wgmma.wait_group. C'est ce qui rend possible le pipelining agressif.

Mesures : wgmma atteint 95 % du pic théorique de Hopper, contre 62,9 % pour l'instruction mma rétro-compatible (microbenchmarking Hopper).

Blackwell : tcgen05, la mémoire dédiée

tcgen05.mma va plus loin : les opérandes vivent en mémoire partagée et l'accumulateur en Tensor Memory (TMEM), une mémoire de 256 Ko par SM entièrement séparée du banc de registres.

Structure de la TMEM : 512 colonnes × 128 lignes de cellules 32 bits, avec 16 To/s en lecture et 8 To/s en écriture par SM (Blackwell GPU Wiki). Elle s'accède avec des instructions dédiées : tcgen05.ld, tcgen05.st, tcgen05.cp.

Deux autres changements majeurs :

  • la paire de CTA : deux blocs sur deux SM voisins peuvent coopérer sur une seule MMA, ce qui porte la tuile maximale à \(256 \times 256 \times 16\) ;
  • la latence quasi constante. Voici les mesures publiées :
Instruction Tuile Latence (cycles)
wgmma (Hopper) m64n64k16 32,0
wgmma (Hopper) m64n128k16 64,0
wgmma (Hopper) m64n256k16 128,0
tcgen05.mma (Blackwell) m64n64k16 11,0
tcgen05.mma (Blackwell) m128n128k16 11,3
tcgen05.mma (Blackwell) m256n256k16 11,4

Source : arXiv:2512.02189. La latence Blackwell est 2,9 à 11,6× plus faible, et surtout elle ne croît pas avec la taille de tuile — ce qui change complètement la façon dont on conçoit un pipeline logiciel.

Débits mesurés sur B200 par précision (tuile m64n8k16) :

Précision TFLOPS Latence (cycles)
FP64 44,8 ~11,2
FP32 481,2 ~11,5
FP16 1 929,2 11,2
FP8 3 851,4 11,8
FP4 7 702,5 12,6

Notez la progression : chaque division par deux de la largeur double le débit, presque exactement. C'est la logique économique qui pousse toute l'industrie vers FP4.

5.4 Écrire du code qui les utilise

Trois niveaux, du plus simple au plus performant.

Niveau 1 — ne rien écrire

Appelez cuBLAS, cuDNN, ou torch.matmul. Ces bibliothèques sont écrites par des gens dont c'est le métier à temps plein, et vous ne les battrez pas sur une GEMM générique.

Niveau 2 — l'API wmma

#include <mma.h>
using namespace nvcuda::wmma;

__global__ void gemm_wmma(const half* A, const half* B, float* C, int N) {
    fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, float>        c_frag;

    fill_fragment(c_frag, 0.0f);

    for (int k = 0; k < N; k += 16) {
        load_matrix_sync(a_frag, A + /* offset */, N);
        load_matrix_sync(b_frag, B + /* offset */, N);
        mma_sync(c_frag, a_frag, b_frag, c_frag);
    }

    store_matrix_sync(C + /* offset */, c_frag, N, mem_row_major);
}

Ce que wmma ne fait pas

L'API wmma est pédagogiquement excellente et plafonne autour de 60 % du pic. Elle ne donne pas accès à wgmma ni à tcgen05, ne gère pas le pipelining asynchrone, et impose des dispositions figées. Utilisez-la pour comprendre, pas pour produire.

Ce code n'a pas été compilé dans l'environnement de rédaction ; il illustre la structure de l'API telle que documentée.

Niveau 3 — CUTLASS, CuTe DSL, Triton, ThunderKittens

C'est là que se fait le vrai travail. Voir Écosystème. Un repère : FlashAttention-4 est écrit entièrement en CuTe DSL et atteint 1 605 TFLOPS sur B200, soit 71 % d'utilisation matérielle (arXiv:2603.05451).

5.5 Le côté AMD : les Matrix Cores

L'équivalent AMD s'appelle Matrix Core et son instruction MFMA (Matrix Fused Multiply-Add) :

\[\mathbf{D} := \mathbf{A} \times \mathbf{B} + \mathbf{C}\]

Elles s'invoquent par des intrinsèques du compilateur :

// CDNA3 (MI300X) : 16x16x16 en FP16
v = __builtin_amdgcn_mfma_f32_16x16x16f16(a, b, c, 0, 0, 0);

CDNA 4 (gfx950, MI350X/MI355X) double le débit des Matrix Cores pour les types ≤ 16 bits et introduit des instructions à échelle par blocs d'exposant (MXFP8/MXFP6/MXFP4) — l'équivalent fonctionnel du block scaling de Blackwell.

Différences structurelles à connaître pour un portage :

NVIDIA AMD CDNA
Portée de l'instruction warp (32), warp group (128), paire de SM wavefront (64)
Opérandes registres, mém. partagée, TMEM registres (VGPR)
Asynchrone oui (wgmma, tcgen05) non (MFMA est synchrone)
Descripteur mémoire oui non

L'absence d'équivalent à wgmma/TMA côté AMD explique en partie pourquoi les megakernels AMD (Kog, Fleet) prennent des chemins différents.

5.6 La sparsité structurée 2:4

Depuis Ampere, les tensor cores savent exploiter une forme de parcimonie très contrainte : dans chaque groupe de 4 poids consécutifs, au moins 2 doivent être nuls. Le matériel stocke alors 2 valeurs + 2 indices sur 2 bits, ce qui divise par deux le volume et double le débit.

Pourquoi c'est peu utilisé

Le motif 2:4 est extrêmement rigide. Élaguer un modèle pour s'y conformer coûte de la qualité, et le ré-entraînement nécessaire n'est pas toujours possible. En pratique, la quantification (FP8, FP4) offre un gain comparable pour bien moins de dégradation, et c'est elle qui a gagné.

Retenez surtout que les chiffres marketing « avec sparsité » sont le double des chiffres réalisables dans le cas général.

Résumé du chapitre

À retenir

  • Un tensor core fait \(\mathbf{D} = \mathbf{A}\mathbf{B} + \mathbf{C}\) sur de petites matrices, en une instruction. Le gain vient de la réutilisation câblée en silicium.
  • Le prix : des formats et des dispositions de registres très contraints. On passe presque toujours par une bibliothèque.
  • Hopper (wgmma) a rendu les opérandes accessibles depuis la mémoire partagée et l'instruction asynchrone. Blackwell (tcgen05) a ajouté une mémoire dédiée (TMEM), la coopération entre deux SM, et une latence quasi constante de ~11 cycles.
  • Chaque division par deux de la précision double le débit : c'est ce qui pousse vers FP8 puis FP4.
  • L'API wmma plafonne à ~60 % du pic ; le vrai travail se fait en CUTLASS/CuTe, Triton ou ThunderKittens.

Vérifiez que vous avez compris

Pourquoi le passage de FP16 à FP8 double-t-il le débit du tensor core, alors que le nombre de multiplieurs ne change pas ?

Parce que ce sont les mêmes multiplieurs, reconfigurés. Un multiplieur 16 bits peut être scindé en deux multiplieurs 8 bits qui travaillent en parallèle — c'est un choix de conception explicite du circuit. Le débit double, la précision est divisée.

Cela dit, le gain réel en bout de chaîne est souvent plus que double, parce que le volume de données à déplacer est aussi divisé par deux. Sur un noyau limité par la mémoire, c'est même le seul gain qui compte.

Un noyau utilise wmma et atteint 55 % du pic BF16 sur H100. Vaut-il la peine de passer à wgmma directement ?

Écrire wgmma à la main est très difficile (descripteurs, dispositions d'accumulateurs, synchronisation asynchrone). La bonne réponse en 2026 est presque toujours : passer par CUTLASS ou CuTe DSL, qui génèrent le wgmma correct et le pipelining associé.

Le gain attendu est réel — wgmma atteint 95 % du pic contre ~63 % pour la mma classique — mais il ne vient pas de l'instruction seule : il vient du pipeline asynchrone (TMA + mbarrier + spécialisation des warps) qu'elle rend possible. Voir CUDA moderne.

Sur B200, la latence de tcgen05.mma ne dépend presque pas de la taille de tuile. En quoi est-ce important pour la conception d'un noyau ?

Sur Hopper, une tuile deux fois plus grande coûtait deux fois plus de cycles : on choisissait la taille de tuile en arbitrant entre réutilisation des données et latence à masquer. Sur Blackwell, la grande tuile est essentiellement gratuite en latence, donc on prend toujours la plus grande qui tient en TMEM et en mémoire partagée.

Le goulot se déplace : ce n'est plus l'instruction MMA qu'il faut masquer, c'est l'alimentation en données. D'où l'importance de TMA et du pipelining, et d'où le co-design de FlashAttention-4, qui déplace le calcul des exponentielles hors des SFU parce que ce sont elles qui deviennent limitantes.

Chapitre suivant : 6 · Les nombres flottants

Sources de ce chapitre