SIMD, GPU et coprocesseurs

Les cœurs du chapitre précédent trouvent du parallélisme entre les instructions : pendant qu’une addition attend, une autre, indépendante, s’exécute. C’est le parallélisme d’instructions, et il plafonne à une poignée d’instructions par cycle. Beaucoup de programmes offrent un parallélisme bien plus simple : la même opération appliquée à des milliers ou des millions de valeurs. Éclaircir une image ajoute le même nombre à chaque pixel. Entraîner un réseau de neurones multiplie de grandes matrices. Ce chapitre porte sur le matériel conçu pour ce parallélisme de données : les instructions SIMD d’un cœur ordinaire, les GPU et d’autres coprocesseurs.

Le parallélisme d’instructions comprend aussi les processeurs VLIW, que traite le chapitre VLIW et Itanium. Un exemple classique de VLIW est le TriMedia de Philips, un processeur multimédia du début des années 2000 qui n’existe plus en tant que gamme ; ses idées survivent dans les processeurs de signal (DSP) comme l’Hexagon de Qualcomm, une conception VLIW présente dans les puces Snapdragon des téléphones. Ce que faisaient les instructions multimédia du TriMedia (travailler sur quatre octets rangés dans un seul registre) se trouve aujourd’hui dans tous les CPU, sous le nom de SIMD.

Une instruction, beaucoup de données

Une instruction SIMD (single instruction, multiple data) traite un registre large comme un petit tableau de voies (lanes) et fait la même opération sur toutes les voies à la fois. Un registre de 128 bits contient seize valeurs de 8 bits, huit de 16 bits, quatre de 32 bits ou deux de 64 bits. Un seul paddd sur x86, ou un seul add v0.4s, v1.4s, v2.4s sur ARM64, additionne quatre paires d’entiers 32 bits dans le temps d’une seule.

On peut saisir l’idée sans matériel spécial. La démo ci-dessous additionne seize octets de deux façons : un par un, puis quatre à la fois, rangés dans des entiers de 32 bits, une astuce appelée SWAR (SIMD within a register). La seule difficulté est la retenue : la retenue d’un octet ne doit pas déborder sur la voie suivante. add4 efface donc le bit de poids fort de chaque voie avant l’addition, ce qui rend impossible toute retenue d’une voie à l’autre, puis remet les bits de poids fort avec un OU exclusif.

Les 16 voies concordent. Regardez la dernière : 200 + 100 donne 44, car chaque voie reboucle modulo 256, exactement comme le paddb du x86. Pour des pixels, passer du clair au sombre par rebouclage est faux ; les jeux d’instructions SIMD ont donc aussi des additions à saturation (paddusb sur x86, uqadd sur ARM64) qui s’arrêtent à 255. Les unités multimédia du TriMedia avaient déjà cette arithmétique saturée.

Dans le simulateur, qui compile sans optimisation, la boucle octet par octet a pris environ 400 instructions et la boucle groupée environ 240, appels de fonction compris. Un vrai matériel SIMD supprime complètement les masques : l’additionneur lui-même est coupé aux frontières des voies, et une seule instruction traite les seize octets. (Le simulateur n’implémente pas SSE, d’où la version faite à la main.)

Les registres SIMD, de SSE à SVE

Tous les grands jeux d’instructions se sont dotés d’extensions SIMD, chacune plus large que la précédente :

L’histoire du x86 est celle d’élargissements successifs, et chaque nouvelle largeur a exigé de nouvelles instructions et des logiciels recompilés. AVX-512 a eu une vie mouvementée : Intel l’a retiré de ses puces de bureau de 12e génération, alors qu’AMD l’a ajouté avec Zen 4. Le M2 sur lequel ce chapitre a été écrit n’a que NEON : sysctl hw.optional liste FEAT_AES, FEAT_SHA256, FEAT_I8MM et FEAT_BF16, mais pas SVE.

SVE chez ARM et l’extension vectorielle de RISC-V adoptent une autre approche : la largeur des registres n’est pas fixée par l’ISA. Chaque puce choisit une largeur, et le code est écrit pour fonctionner avec n’importe laquelle, comme le montre la section suivante.

Une boucle vectorisée

Voici SAXPY (single-precision a·x plus y), un noyau de calcul classique :

void saxpy(int n, float a, const float *restrict x, float *restrict y) {
    for (int i = 0; i < n; i++)
        y[i] = a * x[i] + y[i];
}

Compilé par clang 21 en -O2 pour ARM64, le cœur de la boucle est :

.loop:
    ldp   q1, q4, [x10, #-16]      // load 8 floats of y (two 128-bit registers)
    subs  x12, x12, #8
    ldp   q2, q3, [x11, #-16]      // load 8 floats of x
    add   x11, x11, #32
    fmla  v1.4s, v2.4s, v0.s[0]    // y[0..3] += x[0..3] * a
    fmla  v4.4s, v3.4s, v0.s[0]    // y[4..7] += x[4..7] * a
    stp   q1, q4, [x10, #-16]      // store 8 floats
    add   x10, x10, #32
    b.ne  .loop

Le compilateur l’a fait seul, sans code particulier : c’est l’auto-vectoriseur. Chaque itération traite huit flottants avec deux instructions fmla (multiplication-addition fusionnée) de quatre voies chacune. Le compilateur produit aussi une boucle scalaire ordinaire, un fmadd s1, s0, s1, s2 par élément, pour les derniers éléments quand n n’est pas un multiple de 8.

Les mots-clés restrict comptent. Ils promettent que x et y ne se chevauchent pas. Sans eux, clang doit envisager que y recouvre en partie x, et ajoute un test de chevauchement à l’exécution avant de pouvoir utiliser la boucle vectorielle.

Pour x86-64, la même boucle dépend des extensions que le compilateur a le droit d’utiliser :

Le x86-64 de base ne garantit que SSE2, qui n’a pas de multiplication-addition fusionnée : la compilation par défaut multiplie et additionne séparément. C’est le piège du SIMD x86 : un binaire compilé pour la base n’utilise pas les unités larges de la puce sur laquelle il tourne, sauf si le programme interroge le CPU à l’exécution et choisit un autre chemin de code, comme le font des bibliothèques telles que le memcpy de la glibc.

SVE évite le problème. Compilée pour l’A64FX de Fujitsu (-mcpu=a64fx), la boucle utilise des registres z extensibles et des prédicats :

    ptrue  p0.s                          // predicate: all lanes active
    rdvl   x12, #4                       // x12 = 4 × the vector length, in bytes
.loop:
    ldr    z2, [x13]
    ldr    z6, [x14]
    fmad   z2.s, p0/m, z1.s, z6.s        // multiply-add, on every active lane
    ...
    add    x13, x13, x12                 // advance by however wide the vectors are

Nulle part le code ne dit combien de flottants tient un registre. Il le demande au matériel avec rdvl (read vector length) : le même binaire traite 4 voies à la fois sur une puce SVE de 128 bits et 16 sur l’A64FX.

Ce que le SIMD rapporte sur le M2

Pour le mesurer, le même fichier a été compilé deux fois avec clang -O2 : une fois normalement, une fois avec -fno-vectorize -fno-slp-vectorize, qui désactive le vectoriseur (la version scalaire ne contient plus aucune instruction vectorielle). Chaque noyau a ensuite tourné sur un cœur performance du M2 Ultra, sur un petit tableau qui tient dans le cache L1 (4 096 éléments) et sur un grand qui n’y tient pas (16 millions d’éléments, 64 Mo par tableau). Meilleur de trois essais, répété deux fois avec les mêmes résultats :

Dans le cache, le code vectoriel est cinq à huit fois plus rapide. La somme d’entiers gagne plus que ne le laisseraient attendre quatre voies, parce que le vectoriseur découpe aussi le total en plusieurs accumulateurs indépendants : la boucle scalaire est une seule chaîne dépendante à une addition par cycle, la boucle vectorielle non.

Hors du cache, le gain diminue. Le grand SAXPY déplace 12 octets par élément (lecture de x, lecture de y, écriture de y) ; à 0,13 ns par élément, cela fait environ 90 Go/s pour un seul cœur, et l’arithmétique n’est plus le goulot d’étranglement. Le SIMD rend le calcul bon marché ; il ne rend pas la mémoire plus rapide.

Les GPU : le SIMD à grande échelle

Un GPU pousse l’idée du SIMD aussi loin que possible. Un exemple classique est l’architecture Fermi de NVIDIA (2010), utilisée dans la GeForce GTX 580 : 16 multiprocesseurs de flux (streaming multiprocessors, SM) de 32 « cœurs CUDA » simples chacun, 512 en tout, avec un cache L2 partagé de 768 Ko.

Le modèle de programmation diffère de celui des instructions SIMD. Un programme GPU, un noyau (kernel), est écrit pour un seul thread (« calcule y[i] pour mon i ») et lancé sur des millions de threads. Le matériel regroupe les threads en warps de 32 (terme de NVIDIA ; Apple parle de SIMD groups, AMD de wavefronts). Un warp avance au pas : une instruction est lue et décodée, et les 32 threads l’exécutent, chacun sur ses propres données. NVIDIA appelle cela SIMT, single instruction, multiple threads. C’est du SIMD où chaque voie apparaît au programmeur comme un thread.

Deux conséquences :

• La divergence de branchement. Si les threads d’un même warp prennent des côtés différents d’un if, le warp exécute les deux côtés l’un après l’autre, en désactivant chaque fois les autres voies. Un code plein de branchements dépendant des données gaspille l’essentiel de la machine.
• Masquer la latence par le multithreading. Un GPU n’a ni gros caches ni exécution dans le désordre pour masquer la latence mémoire. À la place, chaque SM garde des dizaines de warps résidents, chacun avec ses propres registres, et passe à un autre warp prêt à chaque cycle où l’un d’eux attend la mémoire. Les moteurs de paquets des processeurs réseau utilisent la même astuce ; le chapitre sur les multicœurs l’examine dans les CPU.

Depuis Fermi, les chiffres ont été multipliés par plus de dix :

Le H100 est une carte pour centre de données, avec moins de matériel graphique et bien plus d’autres choses : un calcul en double précision plus rapide, et des cœurs tensoriels (tensor cores), des unités qui multiplient de petites matrices de nombres de 8 ou 16 bits en une opération, et fournissent bien plus que le chiffre FP32 pour l’apprentissage automatique. Vers 2010, les GPU utilisés pour le calcul général s’appelaient encore « GPGPU » et restaient une niche. Aujourd’hui, l’essentiel des nouveaux supercalculateurs et presque tout l’entraînement d’IA tournent sur des GPU, comme le montre le chapitre sur les grappes et les supercalculateurs.

Un GPU qu’on peut mesurer

Le M2 Ultra intègre un GPU : system_profiler SPDisplaysDataType indique 60 cœurs sur cette machine. Le GPU d’Apple n’est pas une carte séparée : il partage la même mémoire que les cœurs du CPU, et aucune donnée n’a besoin d’être d’abord copiée à travers un bus.

Voici SAXPY sous forme de noyau de calcul Metal, écrit pour un seul thread :

kernel void saxpy(device const float *x [[buffer(0)]],
                  device float *y       [[buffer(1)]],
                  constant float &a     [[buffer(2)]],
                  uint i [[thread_position_in_grid]]) {
    y[i] = a * x[i] + y[i];
}

Lancé sur 67 millions de threads (2²⁶ éléments, 256 Mo par tableau), 20 fois de suite, depuis un petit programme Swift. Metal a indiqué un threadExecutionWidth de 32 pour ce noyau, la taille de ses SIMD groups. Chaque lot de 20 lancements a pris 25,5 ms de temps GPU sur cinq essais : 0,019 ns par élément, soit 630 Go/s de trafic mémoire. Apple annonce 800 Go/s de bande passante mémoire crête pour la puce. Un seul cœur de CPU, avec la version NEON ci-dessus, atteignait environ 90 Go/s. Pour un travail qui ne fait que parcourir la mémoire, le GPU gagne parce qu’il peut garder beaucoup plus de requêtes mémoire en vol.

Les coprocesseurs

Plus généralement, un CPU principal peut confier du travail à des aides spécialisées, appelées coprocesseurs. Trois familles ont évolué différemment.

Les processeurs réseau. Un processeur réseau programmable réunit de nombreux petits moteurs RISC sur une carte, qui traitent les paquets au débit de la ligne pendant que le CPU principal fait autre chose. Vers 2010, l’Ethernet à 40 gigabits était l’étape suivante ; les ports à 400 Gb/s sont aujourd’hui courants dans les centres de données, et ceux à 800 Gb/s sont livrés. Toutes les cartes réseau délèguent désormais au matériel le calcul des sommes de contrôle et le découpage des gros envois en paquets, et les centres de données utilisent des SmartNIC ou DPU (data processing units) qui exécutent les fonctions de réseau, de stockage et de sécurité sur leurs propres cœurs ARM, comme le préfiguraient les processeurs réseau.

La cryptographie. Les coprocesseurs cryptographiques se présentaient autrefois sur des cartes d’extension. Aujourd’hui, le matériel cryptographique le plus répandu est dans le CPU lui-même. Intel a ajouté les instructions AES-NI en 2010 (aesenc effectue une ronde d’AES), et ARMv8 a aese et aesmc ; le hachage SHA a aussi ses instructions. La différence est spectaculaire. OpenSSL 3.6, chiffrant des blocs de 16 Ko en AES-128-CTR sur un cœur du M2, a atteint 13,1 Go/s avec les instructions AES. En forçant le code C portable d’OpenSSL avec OPENSSL_armcap=0, il est tombé à 0,32 Go/s : environ 40 fois plus lent. Du matériel spécialisé pour une tâche courante et figée, c’est le cas idéal pour un coprocesseur, même quand ce « coprocesseur » se réduit à quelques instructions.

Le graphisme et l’apprentissage automatique. Le GPU a commencé comme coprocesseur graphique et il est devenu un processeur parallèle généraliste. La plus récente unité spécialisée est le processeur neuronal (NPU). Le Neural Engine à 32 cœurs du M2 Ultra est annoncé par Apple à 31,6 billions d’opérations par seconde, sur des nombres de faible précision, pour l’inférence de réseaux de neurones. Contrairement au GPU, on ne le programme pas directement : les applications y accèdent via le framework Core ML d’Apple, qui décide de ce qui s’exécute où. Intel, AMD et Qualcomm mettent désormais des NPU similaires dans leurs puces pour portables.

Le même schéma se répète dans tous ces exemples. Une tâche devient assez courante, et assez régulière, pour justifier du matériel dédié. Elle commence sur une puce ou une carte séparée, et finit soit en instructions du CPU (virgule flottante, SIMD, AES), soit en unité placée à côté des cœurs sur la même puce (GPU, NPU).

À retenir

• Les instructions SIMD font la même opération sur chaque voie d’un registre large : 128 bits pour SSE et NEON, 256 pour AVX2, 512 pour AVX-512, et une largeur choisie par la puce pour SVE et les vecteurs de RISC-V.
• Les compilateurs vectorisent automatiquement les boucles simples : clang a transformé SAXPY en fmla v1.4s sur ARM64 et en mulps/addps, vfmadd…ymm ou …zmm sur x86-64 selon les options de la cible.
• Sur un cœur de M2, la vectorisation a rendu les noyaux 5 à 9 fois plus rapides dans le cache, mais seulement 3,5 à 5 fois hors du cache, où la bande passante mémoire (~90 Go/s par cœur ici) est la limite.
• Les GPU exécutent des milliers de threads en warps de 32 qui avancent au pas (SIMT), masquent la latence mémoire en changeant de warp, et souffrent de la divergence de branchement. Le GPU du M2 Ultra a parcouru SAXPY à environ 630 Go/s.
• Les coprocesseurs pour le réseau, la cryptographie, le graphisme et les réseaux de neurones tendent à migrer sur la puce du CPU. Les instructions AES ont rendu le chiffrement environ 40 fois plus rapide que le code portable sur le M2.