Processeurs parallèles et multicœurs

Les ordinateurs parallèles effectuent des opérations en même temps. Ils peuvent réduire la latence, le temps nécessaire à une tâche, ou augmenter le débit, le nombre de tâches achevées par seconde. Ces objectifs conduisent à différentes façons de répartir le travail.

Le mur de puissance

Augmenter la fréquence d’horloge rendait autrefois chaque génération de processeurs plus rapide. La puissance dynamique suit approximativement $P\approx C{V}^{2}f$ : capacité multipliée par la tension au carré et par la fréquence de commutation. La tension intervient deux fois dans ce produit. Réduire la tension et la fréquence de 15 % donne ${\mathrm{0,85}}^2\times \mathrm{0,85}=\mathrm{0,614125}$ de la puissance dynamique initiale, soit environ 61 %.

La tension ne pouvait pas continuer à baisser suffisamment pour compenser l’augmentation des fréquences d’horloge. Les fuites et la chaleur rendaient trop coûteuse la poursuite de leur augmentation. Ce mur de puissance a conduit l’industrie vers les puces multicœurs au milieu des années 2000. Une puce multicœur place plusieurs cœurs de processeur sur une même puce ; un système multiprocesseur possède au moins deux processeurs. Des cœurs plus modestes peuvent accomplir davantage de travail total par joule, à condition qu’il y ait assez de travail à partager.

Quatre formes de travail parallèle

Le matériel et les compilateurs détectent principalement le parallélisme au niveau des instructions. L’exécution en pipeline, l’émission multiple et l’exécution dans le désordre font se chevaucher les instructions tout en préservant les résultats exigés par le programme. Les dépendances et les branchements limitent ce chevauchement.

Le parallélisme de données apparaît lorsque de nombreux éléments nécessitent le même calcul. Les instructions SIMD, les processeurs vectoriels et les GPU exploitent cette régularité. Le parallélisme au niveau des threads exige que le code, une bibliothèque ou un compilateur répartisse le travail entre les threads. Un thread possède son propre compteur ordinal, ses registres et sa pile ; les threads d’un même processus partagent son espace d’adressage. Des processus distincts ont des espaces d’adressage distincts.

Les tâches au niveau des requêtes nécessitent peu de coopération entre elles. Ajouter des serveurs peut donc augmenter la capacité presque proportionnellement aux ressources ajoutées. Exécuter plusieurs programmes indépendants sur un ordinateur portable utilise aussi des cœurs qui seraient autrement inactifs et augmente le débit global, même si aucun de ces programmes n’est parallélisé en interne.

Tâches concurrentes et exécution simultanée

La concurrence signifie que des tâches progressent pendant des périodes qui se chevauchent. Un cœur peut le faire en passant de l’une à l’autre. Le parallélisme signifie que des opérations s’exécutent au même instant sur des unités matérielles distinctes. Un système monocœur peut faire progresser plusieurs programmes sans exécuter leurs instructions simultanément.

Un programme comportant un seul thread séquentiel ne se répartit pas sur huit cœurs simplement parce qu’ils existent. Les cœurs supplémentaires peuvent exécuter d’autres tâches, mais réduire la durée de ce programme nécessite du travail parallèle en son sein. Même après la répartition du travail, toute partie séquentielle restante limite encore le gain.

Un processeur peut contenir de nombreuses unités arithmétiques sans suivre de nombreux flux d’instructions indépendants. La classification de Flynn classe le matériel selon ses flux d’instructions et ses flux de données : la première étape consiste donc à choisir le niveau de la machine décrit.

Les quatre classes matérielles

Un processeur monocœur classique est SISD, même lorsque son pipeline fait se chevaucher plusieurs instructions. L’émission superscalaire ne crée pas à elle seule des compteurs ordinaux indépendants et ne transforme pas le cœur en matériel MIMD.

Une instruction SIMD peut additionner huit paires de nombres à virgule flottante de 32 bits dans un registre de 256 bits. Elle possède huit voies arithmétiques, mais celles-ci suivent une seule instruction. Les extensions SIMD de largeur fixe comprennent x86 SSE et AVX et Arm Neon ; les architectures vectorielles appartiennent aussi à la classe SIMD. Les calculs réguliers sur tableaux conviennent bien, car de nombreux éléments nécessitent la même opération.

Les processeurs MIMD choisissent leurs instructions indépendamment. Les puces multicœurs, les multiprocesseurs à mémoire partagée et les grappes de calcul appartiennent à cette classe. Les cœurs peuvent exécuter des programmes différents ou des copies d’un même programme qui atteignent des branches différentes à des moments différents.

MISD complète la classification à deux par deux. Elle comporte peu d’exemples pratiques incontestés. Certains systèmes tolérants aux pannes et pipelines de traitement sont décrits ainsi, mais leur classification dépend de la façon dont leurs flux sont définis.

Un seul programme peut avoir plusieurs flux d’instructions

SPMD, un programme et plusieurs ensembles de données, est un style de programmation. Chaque unité de travail exécute une copie du même programme et utilise son identifiant pour sélectionner sa part des données. Les unités de travail ont des états d’exécution indépendants : SPMD s’exécute donc couramment sur du matériel MIMD.

Pour quatre cœurs additionnant 4 000 valeurs, le cœur d’identifiant $i$ traite les indices $1000i$ à $1000i+999$. Le cœur 0 traite les indices de 0 à 999 ; le cœur 3 traite ceux de 3000 à 3999. Chacun produit une somme partielle, puis ces quatre sommes sont combinées. Le texte du programme partagé n’exige pas que les quatre cœurs émettent la même instruction au même cycle.

Classer chaque niveau séparément

Considère un processeur à huit cœurs dont les threads utilisent des instructions AVX. Au niveau des cœurs, il est MIMD : les cœurs avancent indépendamment. À l’intérieur d’une opération AVX, il est SIMD : une instruction commande plusieurs voies. Compter ensemble toutes les unités arithmétiques masquerait cette structure.

Une grappe de calcul exécutant des copies d’un programme scientifique est de même un matériel MIMD utilisant un logiciel SPMD. Ses nœuds communiquent par messages. Un GPU combine de nombreux groupes de threads ordonnancés indépendamment avec une exécution de type SIMD à l’intérieur de chaque groupe ; NVIDIA appelle SIMT son mode d’exécution fondé sur les threads.

Pour classer une description, demande-toi qui choisit l’instruction suivante. Demande-toi ensuite si cette instruction traite un ou plusieurs éléments de données. Le nombre de fichiers sources, de programmes installés ou d’unités arithmétiques ne suffit pas à répondre à ces deux questions.

Un défaut de cache peut laisser un cœur en attente pendant des centaines de cycles. Le multithreading matériel conserve les états d’exécution de plusieurs threads dans un même cœur afin que des instructions utiles d’un autre thread puissent s’exécuter pendant cette attente.

Ce qui est dupliqué et ce qui est partagé

Chaque thread matériel possède son propre compteur ordinal et son propre état des registres. Les threads partagent les unités d’exécution, les caches et l’accès à la bande passante mémoire du cœur. Conserver l’état dans le matériel rend le changement bien plus rapide que la sauvegarde et la restauration d’un processus par le système d’exploitation, qui peuvent coûter des milliers de cycles. Le matériel peut changer de thread en un cycle ou moins.

Le système d’exploitation voit chaque thread matériel comme un processeur logique. Huit cœurs physiques avec deux threads matériels par cœur présentent 16 processeurs logiques. Ce nombre décrit les contextes de threads disponibles, et non 16 cœurs physiques complets.

Trois approches d’ordonnancement

L’ordonnancement à grain fin peut passer d’un thread prêt à l’autre même lorsque le thread précédent pourrait continuer. L’ordonnancement à gros grain continue d’exécuter ce thread jusqu’à ce qu’un blocage coûteux justifie un changement. Le SMT utilise un cœur à émission multiple et à ordonnancement dynamique : les instructions de différents threads peuvent occuper différents créneaux d’émission libres pendant un même cycle. Le renommage des registres sépare leurs états. Hyper-Threading d’Intel est une implémentation du SMT ; certains processeurs IBM POWER prennent en charge huit threads matériels par cœur.

Étudier le masquage de la latence

Supposons qu’un thread calcule pendant 20 cycles puis attende la mémoire pendant 200 cycles. Son cycle dure 220 cycles. Seul, son travail utile occupe $20\mathrm{/}220\approx \mathrm{0,091}$, soit 9,1 % du temps.

Avec quatre threads, un travail décalé librement, sans coût de changement ni concurrence supplémentaire, le travail utile peut occuper $4\times 20\mathrm{/}220\approx \mathrm{0,364}$, soit 36,4 %. Le cœur connaît encore des périodes d’inactivité, car quatre séquences de 20 cycles ne peuvent pas couvrir les périodes d’attente.

Onze threads fournissent au total 220 cycles de séquences de calcul. Pendant que l’un attend 200 cycles, les dix autres peuvent calculer chacun pendant 20 cycles. Cet ordonnancement idéal maintient le cœur occupé. La performance réelle dépend aussi de la concurrence éventuelle entre ces threads pour les ressources mémoire et d’exécution.

Débit et durée d’exécution d’un thread

Le multithreading augmente principalement le débit. Un thread donné peut prendre plus de temps, car les autres utilisent des cycles et de l’espace de cache qu’il aurait reçus autrement. Deux threads matériels partageant une unité arithmétique ne peuvent pas tous deux bénéficier de la capacité d’un second cœur complet.

Si les deux threads sollicitent déjà les mêmes unités d’exécution entièrement occupées, le SMT dispose de peu de capacité libre à utiliser. Si l’un est bloqué pendant que l’autre peut s’exécuter, le partage est utile. Les GPU appliquent le masquage de la latence à plus grande échelle en maintenant de nombreux groupes de threads prêts pour leurs ordonnanceurs.

Une machine parallèle est utile lorsqu’elle modifie le temps écoulé. L’accélération compare la meilleure exécution séquentielle à l’exécution parallèle ; l’efficacité compare ce gain au nombre de processeurs utilisés.

Partir du temps mesuré

Pour $p$ processeurs, $S(p)=T_1\mathrm{/}{T}_p$ et $E(p)=S(p)\mathrm{/}p$. Une tâche séquentielle de 100 s qui se termine en 16 s sur huit cœurs a une accélération de $100\mathrm{/}16=\mathrm{6,25}$. Son efficacité est de $\mathrm{6,25}\mathrm{/}8=\mathrm{0,78125}$, soit environ 78 %. Mesure le temps écoulé pour les deux exécutions, avec la même tâche de taille fixe.

Diviser le temps entre parties améliorée et inchangée

Dans la loi d’Amdahl, $f$ est la fraction du temps d’exécution initial qu’une amélioration accélère. C’est une fraction du temps, pas une fraction des lignes de code source. Si cette partie devient $s$ fois plus rapide, alors

La partie inchangée prend toujours la totalité de son temps initial. Pour une répartition parallèle idéale entre $p$ processeurs, pose $s=p$. Ce modèle exclut les surcoûts de communication, de déséquilibre et de synchronisation ; ces coûts peuvent allonger le temps mesuré.

Avec 90 % de temps parallèle et dix processeurs, le temps normalisé est de $\mathrm{0,1}+\mathrm{0,9}\mathrm{/}10=\mathrm{0,19}$. L’accélération est de $1\mathrm{/}\mathrm{0,19}\approx \mathrm{5,26}$ et l’efficacité d’environ 53 %. Avec 100 processeurs, le temps est de $\mathrm{0,1}+\mathrm{0,009}=\mathrm{0,109}$ et l’accélération d’environ 9,17. Multiplier le nombre de processeurs par dix n’a pas doublé l’accélération.

Trouver la limite avant d’acheter plus de processeurs

Lorsque le nombre de processeurs augmente, le temps parallèle tend vers zéro. La fraction séquentielle subsiste et donne le plafond $S_{max}=1\mathrm{/}(1-f)$. Une fraction séquentielle de 10 % limite l’accélération idéale à 10 ; une fraction de 5 % la limite à 20. Avec un nombre fini de processeurs, l’accélération approche ces plafonds par valeurs inférieures.

La loi d’Amdahl permet aussi d’évaluer une classe d’instructions accélérée. Si la multiplication prend 30 % d’une exécution et devient quatre fois plus rapide, le temps normalisé devient $\mathrm{0,7}+\mathrm{0,3}\mathrm{/}4=\mathrm{0,775}$. Le programme entier n’est que $1\mathrm{/}\mathrm{0,775}\approx \mathrm{1,29}$ fois plus rapide.

Remonter à partir d’un objectif

Pour obtenir une accélération de 90 sur 100 processeurs, écris $1\mathrm{/}90=(1-f)+f\mathrm{/}100=1-\mathrm{0,99}f$. On obtient donc $f=(1-1\mathrm{/}90)\mathrm{/}\mathrm{0,99}\approx \mathrm{0,99888}$. Seuls environ 0,11 % du temps initial peuvent rester séquentiels. Un objectif proche du nombre de processeurs exige une très petite fraction inchangée.

Une accélération mesurée peut parfois dépasser le nombre de processeurs. La division d’un grand ensemble de travail peut permettre à chaque partie de tenir dans le cache, réduisant le coût par opération. Cette accélération superlinéaire modifie une hypothèse du modèle simple : le travail parallèle n’est plus simplement le travail initial divisé en parts égales. Vérifie l’algorithme, la taille des données et le comportement du cache avant d’interpréter une telle mesure.

Ajouter des processeurs peut soit raccourcir la même tâche, soit permettre à une machine de traiter une tâche plus grande. Le passage à l’échelle fort et le passage à l’échelle faible mesurent ces utilisations différentes, et tous deux dépendent de la quantité de travail utile qui subsiste après les coûts de coordination.

Travail fixe et travail croissant

Le passage à l’échelle fort maintient la taille du problème fixe. Le passage à l’échelle faible augmente le travail proportionnellement au nombre de processeurs, en maintenant constant le travail par processeur. Un passage à l’échelle faible parfait maintient le temps écoulé constant tandis que le travail total augmente.

Considère dix additions séquentielles suivies d’une somme élément par élément de deux matrices. Une matrice de 10 par 10 nécessite 100 additions parallèles. Le temps total est de 110 unités sur un processeur, 20 sur dix et 12 sur cinquante : les accélérations sont de 5,5 et d’environ 9,2. Une matrice de 40 par 40 nécessite 1 600 additions parallèles. Les temps deviennent 1610, 170 et 42 unités : les accélérations sont d’environ 9,5 et 38,3. Le problème plus grand donne plus de poids au travail parallèle utile par rapport aux mêmes dix additions séquentielles.

La référence différente de Gustafson

La loi de Gustafson utilise la fraction du temps séquentiel mesurée lors de l’exécution parallèle. En appelant cette fraction $s$, l’accélération pour un problème de taille croissante est $S_G=s+(1-s)p$. Pour $s=\mathrm{0,05}$ et 100 processeurs, $S_G=\mathrm{0,05}+\mathrm{0,95}(100)=\mathrm{95,05}$.

La loi d’Amdahl avec 5 % de temps séquentiel lors de l’exécution initiale sur un processeur ne donne qu’environ 16,8 sur 100 processeurs. Ces fractions utilisent des dénominateurs différents et répondent à des questions différentes. Une accélération de 95,05 pour un problème de taille croissante ne signifie pas que la tâche initiale de taille fixe s’est terminée 95,05 fois plus vite.

L’unité de travail la plus lente détermine la fin

Si 400 unités de travail égales sont réparties entre huit cœurs, la part équilibrée est de 50 chacun. Supposons que l’un en reçoive 120 tandis que les sept autres en reçoivent 40 chacun. L’étape dure 120 unités de temps, donc l’accélération est de $400\mathrm{/}120\approx \mathrm{3,33}$. Diviser par la charge moyenne ne tiendrait pas compte de l’attente du cœur le plus chargé.

La création et l’ordonnancement des threads, les messages, les barrières et la concurrence pour la bande passante partagée consomment aussi du temps. La granularité est la quantité de calcul entre deux événements de communication ou de synchronisation. De plus gros morceaux réduisent le coût de coordination par opération utile, mais un nombre insuffisant de morceaux peut laisser certaines unités de travail inactives.

Davantage de processeurs peuvent augmenter le temps

Pour une tâche entièrement parallèle de 100 s avec un surcoût de $\mathrm{0,5}p$ secondes, $T_p=100\mathrm{/}p+\mathrm{0,5}p$. Dix processeurs prennent 15 s. Quatorze prennent environ 14,14 s, près du minimum situé autour de $\sqrt{200}$ processeurs. Quarante prennent 22,5 s, car leur surcoût de 20 s dépasse le gain obtenu par un calcul utile plus court.

Lorsque tu présentes un résultat, précise si le travail est resté fixe ou a augmenté, et compare avec le temps du meilleur programme séquentiel. Utiliser comme référence une implémentation parallèle lente exécutée sur un seul processeur peut gonfler l’accélération annoncée.

Le débit arithmétique annoncé pour un processeur n’est qu’une limite parmi d’autres. Le modèle Roofline permet aussi de déterminer si la mémoire peut fournir les données assez vite pour soutenir ce débit. Il donne une borne supérieure qui aide à choisir une piste d’optimisation utile.

Compter les opérations par octet transféré depuis ou vers la DRAM

L’intensité arithmétique est le nombre d’opérations à virgule flottante divisé par le nombre d’octets transférés entre la DRAM et le processeur. Les travaux d’origine sur le modèle Roofline appellent cette mesure fondée sur la DRAM intensité opérationnelle. Compte les transferts à cette frontière, en incluant les lectures et les écritures.

DAXPY calcule $y_i=a{x}_i+y_i$ avec des nombres en double précision. Chaque élément nécessite une multiplication et une addition : deux opérations à virgule flottante. Lire une valeur de 8 octets dans chaque tableau et écrire 8 octets en retour transfère 24 octets. L’intensité est de $I=2\mathrm{/}24\approx \mathrm{0,0833}$ opération à virgule flottante par octet. Le coefficient scalaire est réutilisé ; dans ce modèle de trafic, il n’est pas chargé à nouveau depuis la DRAM pour chaque élément.

Appliquer les deux limites

Si le débit de calcul de crête est $P$ et la bande passante mémoire de crête $B$, le toit est $\min (P,BI)$. La bande passante multipliée par le nombre d’opérations à virgule flottante par octet s’exprime en opérations à virgule flottante par seconde. Sur un graphique à échelles logarithmiques, la borne de bande passante monte et la borne de calcul est horizontale.

Pour 1 000 GFLOP/s et 100 Go/s, le point de transition est de $P\mathrm{/}B=1000\mathrm{/}100=10$ opérations à virgule flottante par octet. En dessous de cette intensité, le toit de bande passante est plus bas ; au-dessus, le toit de calcul est plus bas. DAXPY est limité à $100(2\mathrm{/}24)\approx \mathrm{8,33}$ GFLOP/s, soit moins de 1 % du débit de calcul de crête de la machine. Un noyau de calcul d’intensité 2 est limité à 200 GFLOP/s.

Un noyau limité par la mémoire peut encore bénéficier de cœurs supplémentaires ou d’unités SIMD plus larges tant qu’il a trop peu de requêtes en cours pour utiliser la bande passante disponible. Une fois cette bande passante saturée, une capacité arithmétique supplémentaire ne relève pas son toit. Augmenter la bande passante ou effectuer plus de calcul utile par octet transféré le relève.

La réutilisation modifie l’intensité

Multiplier deux matrices de 1000 par 1000 en double précision effectue environ $2(1000{)}^3=2\times {10}^9$ opérations à virgule flottante. Si chaque élément des deux entrées et de la sortie traverse une fois la frontière de la DRAM, le trafic total est de 24 Mo, soit environ 83 opérations à virgule flottante par octet. C’est au-dessus du point de transition de la machine de l’exemple.

Une implémentation naïve peut relire plusieurs fois les valeurs des matrices depuis la DRAM. Le traitement par blocs organise les boucles autour de sous-matrices qui tiennent dans le cache, de sorte que les valeurs servent à de nombreuses opérations avant d’être remplacées. L’algorithme reste une multiplication matricielle, mais l’implémentation transfère moins d’octets depuis ou vers la DRAM par opération.

Un toit n’est pas un résultat mesuré

Un noyau peut s’exécuter sous la borne à cause de dépendances, du surcoût des instructions, de mauvaises séquences d’accès ou d’un manque de travail parallèle. Des plafonds plus bas montrent ces restrictions : l’absence de SIMD peut abaisser le plafond de calcul, tandis que de mauvaises séquences d’accès réduisent la bande passante utilisable. Un point sous un toit mérite d’être étudié ; le toit seul ne prédit pas son débit d’exécution exact.

Une extension SIMD divise un registre large en voies. Une instruction effectue la même opération dans chaque voie, ce qui convient bien aux boucles régulières sur des éléments indépendants.

Compter les voies et les itérations

Le nombre de voies est la largeur du registre divisée par celle d’un élément. Un registre de 256 bits contient huit nombres à virgule flottante de 32 bits, quatre nombres en double précision de 64 bits ou 32 éléments d’un octet. SSE utilise des registres XMM de 128 bits ; AVX et AVX2 utilisent des registres YMM de 256 bits. En mode 64 bits, il y a 16 registres YMM. AVX-512 fournit 32 512 bits de registres ZMM, contenant chacun 16 nombres à virgule flottante. Arm Neon utilise des registres de 128 bits.

Pour c[i] = a[i] + b[i] sur 1 000 nombres à virgule flottante, des instructions à huit voies traitent le tableau en 125 itérations vectorielles. Chaque itération charge deux vecteurs, additionne leurs voies et range un vecteur. Avec 1 003 éléments, 125 itérations complètes laissent trois éléments. Une boucle scalaire pour le reste ou une opération masquée traite cette courte fin sans accéder au-delà des tableaux.

Huit voies représentent une capacité arithmétique, pas une garantie que le programme sera huit fois plus rapide. La préparation, le traitement du reste, le travail séquentiel et les accès soumis à la bande passante mémoire prennent toujours du temps.

Choisir une organisation qui remplit un registre

Les chargements vectoriels simples fonctionnent bien sur des éléments contigus, avec un pas de 1. Un tableau de points {x, y, z} dont chaque champ occupe 4 octets place les champs x consécutifs à 12 octets d’intervalle. Une structure de tableaux place ensemble les valeurs x, puis les valeurs y, puis les valeurs z. Traiter uniquement x peut alors remplir un vecteur avec un seul chargement contigu. Les instructions de rassemblement et de dispersion prennent en charge les adresses dispersées, mais généralement avec un coût d’accès supérieur.

Traiter les conditions dans les voies

Pour c[i] = (a[i] > 0) ? a[i] : 0, une comparaison crée un masque par voie. Une sélection choisit entre les résultats possibles, ou des instructions masquées ne mettent à jour que les voies sélectionnées. Les voies inactives conservent ou reçoivent des valeurs selon le comportement du masque de l’instruction.

Si différentes voies nécessitent des branches longues et différentes, les instructions de ces chemins doivent tout de même être émises, avec seulement les voies concernées actives. Le groupe ne peut pas donner à chaque voie un flux d’instructions distinct. Un masque préserve la correction, mais ne peut pas faire disparaître toutes les instructions de ces chemins.

Vérifier l’indépendance des itérations

Les compilateurs peuvent vectoriser automatiquement les boucles ; les programmeurs peuvent aussi utiliser des bibliothèques ou des fonctions intrinsèques comme _mm256_add_ps. Le compilateur doit établir que des itérations simultanées préservent le résultat. La récurrence a[i] = a[i-1] + b[i] dépend de la nouvelle valeur produite par l’itération précédente et ne peut pas être directement vectorisée comme des éléments indépendants. Un chevauchement possible des pointeurs, appelé aliasing, peut aussi nécessiter une vérification du chevauchement à l’exécution ou empêcher la vectorisation.

Une réduction conserve plusieurs sommes partielles dans un registre, puis combine les voies. Modifier l’ordre des additions en virgule flottante peut changer les derniers bits, car l’addition n’est pas associative. Les sommes d’entiers concordent lorsque l’arithmétique choisie ne présente pas de dépassement de capacité ou utilise une arithmétique modulaire définie.

SIMD et l’exécution multicœur peuvent être combinés. Huit cœurs émettant chacun une addition sur huit voies par cycle peuvent effectuer jusqu’à 64 additions à virgule flottante par cycle. Ils restent soumis à des limites partagées comme la bande passante mémoire.

Une architecture vectorielle exprime une opération sur une suite d’éléments et stocke la longueur vectorielle active dans un registre. Le même programme peut s’exécuter sur des implémentations dont les registres vectoriels ont des capacités différentes.

Distinguer capacité, longueur et voies

L’extension vectorielle RISC-V, RVV, version 1.0, a été ratifiée en 2021. Elle fournit 32 registres vectoriels, de v0 à v31. Chaque implémentation choisit la largeur de ses registres, VLEN, sous forme d’une puissance de deux. La largeur des éléments et le regroupement des registres déterminent combien d’éléments peuvent y tenir.

La longueur active, vl, est le nombre d’éléments traités par les prochaines instructions vectorielles. vsetvli configure l’opération vectorielle à partir du nombre d’éléments restants et du format d’élément souhaité, puis renvoie la valeur vl sélectionnée. Le code avance selon la valeur renvoyée. C’est important, car un code portable doit utiliser la longueur effectivement sélectionnée par l’implémentation.

Les voies sont les pipelines parallèles qui exécutent l’opération vectorielle. Une addition de 32 éléments sur quatre voies prend environ huit cycles après le démarrage du pipeline. Un vecteur de 32 éléments ne nécessite pas 32 unités arithmétiques. Des vecteurs plus longs et davantage de voies sont des choix de conception distincts.

Découper une boucle en bandes

Le découpage en bandes divise une longue boucle en morceaux de taille adaptée aux vecteurs. Pour 100 éléments et un maximum de 32 par opération, une suite de longueurs possible est 32, 32, 32 et 4. Chaque passage charge la partie suivante, la traite, la range et réduit le nombre d’éléments restants de sa valeur vl.

DAXPY sur 1 000 nombres en double précision peut de même utiliser 31 passages de 32 suivis d’un passage de huit, soit 32 passages au total. RVV peut choisir d’autres longueurs autorisées vers la fin : la boucle utilise donc la valeur vl renvoyée au lieu de supposer une répartition finale particulière. Définir la longueur active permet de traiter la courte fin sans boucle scalaire distincte pour le reste.

Le schéma de contrôle de la boucle est le suivant :

1. Demander une longueur pour les éléments restants.
2. Charger et traiter ce nombre d’éléments.
3. Ranger les résultats.
4. Faire avancer les pointeurs et soustraire la longueur renvoyée.

Une instruction vectorielle décrit de nombreuses opérations sur les éléments, ce qui réduit le travail de chargement et de décodage par rapport à un corps de boucle scalaire répété pour chaque élément.

Alimenter les opérations dépendantes et traiter les données non contiguës

Le chaînage permet à une instruction vectorielle qui utilise un résultat de commencer à traiter les premiers éléments avant que l’instruction qui les produit ait terminé les suivants. Pour une multiplication suivie d’une addition, le premier produit peut entrer dans le pipeline d’addition pendant que les produits suivants sont encore en cours de calcul.

Un masque, souvent dans le registre RVV v0, sélectionne les éléments à mettre à jour. Un chargement à pas constant accède à des éléments séparés par un nombre fixe d’octets. Dans une matrice de 100 par 100 nombres en double précision rangée par lignes, descendre d’une ligne dans une colonne fait avancer de 100 éléments, soit 800 octets. Les chargements et rangements indexés utilisent des adresses explicites pour les données irrégulières.

Les architectures vectorielles et les extensions de largeur fixe sont toutes deux SIMD dans la classification de Flynn. Les extensions de largeur fixe codent des largeurs particulières ; le code RVV exprime des longueurs qui s’adaptent au matériel disponible. Les deux nécessitent toujours un ordre d’exécution qui respecte les dépendances de données de la boucle.

Un processeur concentre ses ressources sur une faible latence pour quelques threads. Un GPU concentre ses ressources sur un débit élevé pour de nombreux threads parallèles au niveau des données. Choisir entre les deux exige d’examiner l’ensemble de la tâche, y compris les transferts de données.

Les threads s’exécutent sur des multiprocesseurs de flux

Un GPU NVIDIA contient de nombreux multiprocesseurs de flux, ou SM. Un SM possède des voies arithmétiques, des ordonnanceurs de warps, un grand banc de registres et une mémoire partagée rapide sur la puce. Un « cœur CUDA » désigne une voie arithmétique, et non un cœur complet et indépendant comparable à celui d’un processeur. AMD appelle le groupe de traitement correspondant une unité de calcul.

NVIDIA regroupe les threads en warps de 32. Une instruction émise agit sur les threads actifs du warp, chacun ayant ses propres valeurs de registre et données. C’est le SIMT : une instruction, plusieurs threads. AMD utilise des wavefronts de 32 ou 64 threads selon l’architecture.

Un SM conserve de nombreux warps résidents. Lorsque l’un attend la mémoire, l’ordonnanceur sélectionne un autre warp prêt. Disposer de plus de threads prêts que de voies masque la latence ; cela ne réduit pas la latence d’un accès individuel à la DRAM.

La divergence laisse des voies inactives

Si les threads d’un warp suivent des branches différentes, le warp exécute les chemins nécessaires en masquant les autres threads. Une structure conditionnelle dont les threads se répartissent également entre les deux branches, de coût égal, peut prendre environ deux fois le temps d’un chemin, car les deux doivent s’exécuter. Un branchement peut être sans conséquence lorsque sa condition est uniforme dans chaque warp, même si différents warps choisissent des chemins différents.

Coordonner les accès mémoire

Les registres propres à chaque thread contiennent des valeurs privées. La mémoire partagée propre à chaque bloc est un espace sur la puce géré par le programme. Les caches L1 et L2 permettent la réutilisation automatique, tandis que la mémoire globale est la DRAM du GPU, par exemple GDDR ou HBM.

Lorsque 32 threads lisent des mots consécutifs de 4 octets, leur plage de 128 octets peut être servie par un petit nombre de transactions larges. Cette coalescence des accès mémoire utilise efficacement la bande passante. Des adresses dispersées peuvent nécessiter de nombreuses transactions pour les mêmes données utiles. La réutilisation en mémoire partagée peut encore réduire les transferts répétés avec la mémoire globale.

Inclure le temps de transfert

Pour un GPU dédié, les entrées et les résultats traversent une liaison telle que PCIe ou NVLink. Envoyer 800 Mo et recevoir 400 Mo à un débit effectif de 20 Go/s transfère 1,2 Go en 0,06 s, soit 60 ms. Ajouter un noyau de calcul de 10 ms donne 70 ms au total. Un processeur qui termine en 50 ms l’emporte dans cette comparaison malgré le calcul bien plus rapide du GPU.

Conserver les données sur le GPU pour plusieurs noyaux évite des transferts répétés. De grandes tâches régulières peuvent aussi fournir assez d’opérations arithmétiques utiles pour compenser les coûts de transfert et de lancement. Les petites tâches séquentielles riches en branchements favorisent souvent un processeur.

Se spécialiser davantage lorsque le travail le permet

Un système hétérogène utilise le processeur pour exécuter le programme et le système d’exploitation, et un accélérateur pour les parties adaptées. Les accélérateurs spécialisés tels que les TPU et les NPU consacrent du matériel à une tâche plus restreinte, souvent la multiplication matricielle avec des entiers de 8 bits ou des nombres à virgule flottante de 16 bits. La spécialisation peut améliorer la performance par watt dans ce domaine, tout en réduisant l’éventail des travaux que le matériel traite efficacement.

Un noyau de calcul pour GPU s’exécute une fois par thread lancé. L’hôte lance une grille de blocs de threads et chaque thread calcule la partie des données qui lui est affectée. CUDA appelle le processeur l’hôte et le GPU le périphérique ; OpenCL utilise les notions associées d’élément de travail et de groupe de travail.

Calculer un indice unique

Pour un lancement à une dimension, l’indice de l’élément est blockIdx.x * blockDim.x + threadIdx.x. Avec 256 threads par bloc, le thread 5 du bloc 2 traite l’élément 517 : les deux blocs précédents représentent 512 éléments, et le thread ajoute un déplacement de cinq. Les indices des threads et des blocs commencent à zéro.

Un noyau SAXPY applique une multiplication et une addition à chaque nombre à virgule flottante :

__global__ void saxpy(int n, float a, float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];
    }
}

La grille remplace la boucle ordinaire sur les éléments. Chaque thread actif met à jour son propre y[i] ; les threads voisins utilisent des adresses voisines, ce qui favorise la coalescence.

Arrondir le nombre de blocs à l’entier supérieur

Pour 1 000 000 d’éléments et 256 threads par bloc, la grille nécessite $\lceil 1000000\mathrm{/}256\rceil =3907$ blocs. Ceux-ci lancent 1 000 192 threads. Les 192 derniers ont des indices hors du tableau, donc la condition i < n empêche leurs accès mémoire. Arrondir à l’entier inférieur lancerait trop peu de threads et laisserait des éléments valides non traités.

Un bloc de 256 threads forme huit warps. Les blocs peuvent être affectés à n’importe quel SM et s’exécuter dans n’importe quel ordre. La correction du programme ne doit pas dépendre du fait que le bloc 0 se termine avant le bloc 1.

Organiser le travail de l’hôte et du périphérique

Un hôte alloue généralement l’espace de stockage du périphérique, copie les entrées, lance saxpy<<<3907, 256>>>(n, 2.0f, d_x, d_y) et récupère les résultats. Le lancement est asynchrone : le processeur peut continuer pendant l’exécution du noyau. L’hôte doit attendre par une synchronisation ou une copie bloquante adaptée des résultats avant d’utiliser les données calculées.

Coopérer à l’intérieur d’un bloc

Les threads d’un bloc peuvent échanger des valeurs par la mémoire partagée et se retrouver à __syncthreads(). Chaque thread participant à cette barrière du bloc doit l’atteindre. Placer la barrière dans if (i < n) n’est pas sûr lorsque seuls certains threads satisfont la condition.

Pour une réduction de 1 024 valeurs en mémoire partagée, 512 threads additionnent des paires, puis 256 combinent les résultats, puis 128, et ainsi de suite. Dix étapes divisant le nombre de participants par deux laissent une somme. Une barrière entre les étapes empêche un thread d’utiliser un résultat partiel avant que le thread qui le produit l’ait écrit.

Un lancement ordinaire ne comporte pas de barrière à l’échelle de la grille. Des blocs distincts peuvent écrire des sommes partielles qu’un second noyau combine, ou des opérations atomiques adaptées peuvent mettre à jour un résultat global. Une barrière de bloc ne peut pas à elle seule coordonner différents blocs. Une boucle dont le pas est lié à la taille de la grille constitue une autre façon de répartir le travail : chaque thread traite plusieurs éléments en avançant selon la taille totale de la grille.

Un multiprocesseur à mémoire partagée donne à ses processeurs un espace d’adressage physique unique. Les threads communiquent par des chargements et des rangements dans des variables partagées, tout en conservant des valeurs privées telles que des sommes partielles locales.

Les adresses partagées peuvent avoir des coûts différents

En UMA, accès mémoire uniforme, le temps d’accès à la mémoire est à peu près le même depuis chaque processeur. En NUMA, accès mémoire non uniforme, un socket accède plus vite à sa mémoire locale qu’à celle d’un autre socket. L’espace d’adressage est partagé même si la mémoire physique est distribuée.

Avec 70 % d’accès locaux à 80 ns et 30 % d’accès distants à 140 ns, la moyenne pondérée est de $\mathrm{0,7}(80)+\mathrm{0,3}(140)=98$ ns. Déplacer les données près de leur utilisateur principal réduit la part des accès distants. Compter chaque accès comme local sous-estimerait le temps.

Une organisation courante donne à chaque cœur des caches L1 privés pour les instructions et les données, souvent un cache L2 privé, et un cache de dernier niveau partagé. Un anneau ou un maillage relie les cœurs et les caches. La mémoire partagée décrit l’adressage ; elle n’exige pas que chaque cœur utilise un seul cache commun.

Répartir, calculer et réduire

Pour additionner 64 000 valeurs sur 64 cœurs, chaque cœur traite 1 000 valeurs. Les 64 sommes partielles peuvent être combinées dans un arbre de réduction. Le nombre d’additionneurs actifs est divisé par deux à chaque étape : 32, 16, 8, 4, 2 et 1. Six étapes combinent les résultats partiels.

Avec environ 1 000 étapes d’addition locales et six étapes de combinaison, on obtient environ 1 006 étapes sur le chemin le plus long, avant les surcoûts. L’arbre évite que chaque entrée mette à jour un compteur partagé unique. Chaque étape dépend des valeurs achevées de l’étape précédente ; les unités de travail se retrouvent donc à des barrières de synchronisation entre les étapes.

OpenMP peut exprimer cela avec #pragma omp parallel for reduction(+:sum). Cette directive répartit les itérations, donne aux threads des sommes partielles privées et les combine. Les threads POSIX offrent un contrôle plus explicite de la gestion des threads.

Rendre une mise à jour indivisible

Une opération de lecture-modification-écriture atomique rend une mise à jour indivisible du point de vue des autres processeurs. L’échange atomique et la lecture et addition atomiques en sont des exemples. Un chargement suivi d’un rangement ordinaire représente deux opérations distinctes, entre lesquelles un autre thread peut intervenir.

RISC-V fournit lr.w et sc.w. Le chargement avec réservation lit un mot et établit une réservation. Le rangement conditionnel tente l’écriture et indique sa réussite ou son échec. Une écriture entre-temps à l’emplacement réservé fait échouer la tentative ; le logiciel réessaie. La réservation ne bloque pas les autres processeurs, et l’échec peut aussi survenir pour d’autres raisons autorisées. Ces primitives servent aux mécanismes de synchronisation.

Surveiller les goulots d’étranglement partagés

Si chaque cœur peut lire en continu à 15 Go/s mais que le socket fournit 60 Go/s, environ quatre flux de ce type saturent la bande passante. Davantage de cœurs ne peuvent pas augmenter le débit de données du socket. La communication par les caches a aussi un coût : un cœur qui écrit sans permission d’écriture doit obtenir la propriété et invalider les autres copies. Un cœur qui détient déjà un bloc à l’état modifié peut continuer localement ; les lecteurs ultérieurs peuvent devoir charger une nouvelle copie.

Les caches privés améliorent la latence, mais peuvent détenir plusieurs copies d’un même emplacement mémoire. La cohérence des caches empêche ces copies de donner des historiques contradictoires des écritures à cet emplacement.

Pourquoi la mise à jour de la mémoire ne suffit pas

Partons de X = 0 en mémoire. Le cœur A lit et conserve zéro dans son cache ; le cœur B fait de même. A écrit un dans son cache à écriture différée. Sans protocole de cohérence, B peut trouver sa propre copie en cache et continuer à lire zéro. Une écriture simultanée en mémoire ne corrigerait pas la copie déjà présente dans le cache de B.

La cohérence fonctionne à la granularité du bloc de cache, généralement de 64 octets. Un contrôleur suit l’état de chaque bloc et applique les règles du protocole lorsque les processeurs le lisent ou l’écrivent. La migration déplace les données utiles dans le cache de leur utilisateur actuel. La réplication permet à plusieurs lecteurs de détenir des copies non modifiées. La cohérence doit préserver ces deux avantages tout en contrôlant les écritures.

Trois propriétés requises

Si un emplacement reçoit une écriture de un suivie d’une écriture de deux, les lecteurs ne peuvent pas être en désaccord sur l’écriture la plus récente. Observer l’écriture ultérieure puis l’écriture antérieure inverserait cet historique. Une nouvelle écriture ultérieure pourrait bien sûr stocker à nouveau la valeur un ; la sérialisation concerne les événements d’écriture, pas la possibilité qu’une valeur numérique réapparaisse.

Contrôler la propriété et les copies

Dans un protocole d’invalidation à l’écriture, un cœur qui écrit obtient une permission d’écriture exclusive et invalide les autres copies en cache. Une lecture ultérieure d’un bloc invalidé produit un défaut et charge les données actuelles. C’est un défaut de cohérence, qui s’ajoute aux défauts à froid, de capacité et de conflit.

Un protocole de mise à jour à l’écriture envoie les nouvelles valeurs aux copies existantes. Des écritures répétées peuvent rendre cette méthode coûteuse. L’espionnage de bus diffuse les requêtes de cohérence afin que les caches puissent réagir. Un répertoire suit les détenteurs des copies et dirige les messages vers les caches concernés.

Cohérence, ordre et atomicité

La cohérence concerne un emplacement. Un modèle de consistance mémoire définit les ordres observables autorisés et les résultats des lectures et écritures, y compris les opérations sur différents emplacements. La consistance séquentielle exige un résultat correspondant à un entrelacement qui préserve l’ordre du programme de chaque processeur. Les modèles plus faibles autorisent davantage de réordonnancements ; les instructions de barrière mémoire imposent l’ordre requis lorsque c’est nécessaire.

Ni des copies cohérentes ni un ordre conforme au modèle de consistance ne rendent atomique une incrémentation composée de plusieurs instructions. Supposons que deux threads incrémentent un compteur partagé initialement égal à cinq. Une séquence possible est : A lit cinq, B lit cinq, A écrit six, B écrit six. Chaque lecture et chaque écriture peuvent être cohérentes, mais le compteur final vaut six au lieu de sept. Le programme a besoin d’une mise à jour atomique ou d’une synchronisation adaptée pour préserver les deux incrémentations.

Un contrôleur de cache utilisant l’espionnage surveille les requêtes de cohérence sur une interconnexion à diffusion. Pour suivre un protocole, conserve un état pour la copie du bloc de chaque cœur et mets à jour chaque copie concernée après chaque accès.

Lire les états comme des permissions

Dans MSI de base, un défaut de lecture envoie une requête de lecture sur le bus et charge le bloc à l’état S. Écrire dans un bloc à l’état S nécessite une demande de passage en exclusif pour invalider les autres copies. Écrire depuis l’état I nécessite une requête de lecture exclusive pour charger les données et obtenir la propriété. Le cœur qui écrit termine à l’état M. Un succès à l’état M ne nécessite aucune transaction de cohérence.

Suivre cinq accès

Partons de X = 0 en mémoire et de A et B tous deux à l’état invalide.

1. A lit X. Une lecture sur le bus fournit zéro. États : A = S, B = I.
2. B lit X. Une autre lecture sur le bus fournit zéro. États : A = S, B = S.
3. A écrit un. Sa demande de passage en exclusif invalide B. États : A = M, B = I.
4. B lit X. B subit un défaut ; A fournit un et rétrograde à l’état S. Dans MSI de base, la mémoire est aussi mise à jour. États : A = S, B = S.
5. B écrit deux. Sa demande de passage en exclusif invalide A. États : A = I, B = M.

Chaque accès de cet exemple nécessite une transaction sur le bus. Le quatrième accès est la vérification importante de la propriété : la mémoire seule contenait encore l’ancienne valeur avant que A fournisse ses données modifiées.

Ce que l’état exclusif économise

MESI ajoute E pour une copie unique en cache non modifiée. Si A lit un bloc absent des caches et qu’aucun autre cache ne le détient, A reçoit l’état E. Une écriture ultérieure fait passer E à M sans transaction, car aucune autre copie ne doit être invalidée. MSI chargerait d’abord le bloc à l’état S, puis enverrait une demande de passage en exclusif pour cette écriture.

Si un autre cœur lit le bloc pendant que A le détient à l’état E, les deux passent à l’état S. M et E ont tous deux un seul détenteur valide, mais seul M contient des données plus récentes que la mémoire.

Écritures répétées et passage à l’échelle

Un cœur qui a obtenu l’état M peut mettre à jour le bloc dix fois sans dix nouvelles invalidations. Une méthode de mise à jour à l’écriture enverrait les valeurs modifiées aux détenteurs des copies.

L’espionnage par diffusion devient coûteux lorsque davantage de cœurs surveillent chaque requête. Les structures à répertoire réduisent cette diffusion. Des protocoles plus élaborés ajoutent des états utiles : l’état propriétaire de MOESI permet une copie modifiée partagée, et MESIF désigne un détenteur chargé de transmettre les données en réponse aux requêtes.

Diffuser chaque requête de cohérence devient coûteux sur les grandes machines. Un répertoire enregistre les caches susceptibles de détenir chaque bloc afin que les requêtes puissent leur être envoyées. Ce suivi au niveau du bloc explique aussi un problème courant de performance : le faux partage.

Interroger le nœud d’attache d’un bloc

Chaque bloc possède une entrée dans le répertoire de son nœud d’attache, qui indique s’il est absent des caches, partagé ou modifié et quels caches le détiennent. Les nœuds d’attache sont répartis entre les nœuds ou les parties du répertoire, ce qui distribue le stockage et le traitement des requêtes.

Un défaut de lecture est envoyé au nœud d’attache. Si la mémoire contient les données actuelles, elle peut les fournir et le répertoire enregistre le nouveau détenteur. Si un autre cache possède une copie à l’état modifié, le nœud d’attache lui transmet la requête afin qu’il fournisse les données actuelles. Dans un protocole à répertoire de base, cela remet aussi la mémoire à jour.

Pour écrire sans déjà détenir le bloc à l’état modifié, un cœur demande une permission au nœud d’attache. Le répertoire envoie des invalidations aux détenteurs enregistrés, attend les accusés de réception et accorde la propriété exclusive. Les succès d’écriture répétés à l’état M ne nécessitent ensuite aucun autre message de cohérence. Le passage par le répertoire ajoute de la latence, mais le trafic point à point évite d’interroger chaque cache au sujet de chaque bloc. Les filtres d’espionnage sur puce remplissent une fonction apparentée en limitant les requêtes aux détenteurs possibles.

Estimer le stockage du répertoire

Un vecteur de présence complet comporte un bit par cœur et par bloc. Pour 64 cœurs, cela représente 64 bits, soit huit octets. Par rapport à un bloc de données de 64 octets, les seuls bits de présence coûtent $8\mathrm{/}64=\mathrm{0,125}$, soit 12,5 %. Avec 256 cœurs, le vecteur de 32 octets coûte 50 %. Les bits d’état s’y ajoutent. Des listes limitées de détenteurs, des vecteurs moins détaillés et le suivi des seuls blocs en cache réduisent ce besoin de stockage.

Reconnaître le faux partage

Le vrai partage se produit lorsque des cœurs échangent des données par le même mot. Le faux partage se produit lorsqu’ils accèdent à des variables différentes dans un même bloc de cache et qu’au moins l’un d’eux écrit. Les variables sont indépendantes dans le programme, mais l’unité de cohérence est plus grande que chacune d’elles.

Deux threads incrémentent des compteurs distincts de 8 octets aux adresses 0x1000 et 0x1008. Avec des blocs de 64 octets, leurs numéros de bloc sont $\lfloor 4096\mathrm{/}64\rfloor =64$ et $\lfloor 4104\mathrm{/}64\rfloor =64$. Chaque changement de propriété transfère ensemble les droits sur les deux compteurs. Des écritures alternées fréquentes peuvent faire passer le bloc sans cesse d’un cache à l’autre, même si aucun thread n’a besoin de la valeur de l’autre.

Déplacer le second compteur à 0x1040 le place dans le bloc 65. Le remplissage ou l’alignement peut séparer les données propres à chaque thread ; regrouper les valeurs selon le thread qui les écrit aide aussi. Un accumulateur local réécrit une seule fois réduit le trafic lié à la propriété sans ajouter de remplissage à chaque valeur.

Un cœur qui écrit et un autre qui lit une variable voisine peuvent aussi subir du faux partage : le premier invalide le bloc du lecteur, qui serait autrement utile. Deux lecteurs seuls ne le provoquent pas. Le faux partage ajoute du trafic et de l’attente ; il ne rend pas à lui seul incorrectes les valeurs produites par des mises à jour de variables distinctes. Mesure les accès et le placement des lignes de cache avant de modifier l’organisation des données.

Une grappe de calcul relie des ordinateurs complets possédant chacun leur propre mémoire et leur propre système d’exploitation. Les nœuds coopèrent par des messages explicites. Le placement des données et le coût de communication font donc partie du modèle de performance du programme.

Répartir les données et combiner les résultats

MPI, l’interface de passage de messages, donne à chaque processus un rang et des opérations telles que l’envoi, la réception, la diffusion et la réduction. Un programme SPMD peut utiliser les rangs pour sélectionner les données de chaque processus. Un pointeur dans l’espace d’adressage privé d’un nœud n’est pas une référence ordinaire en mémoire partagée vers les données d’un autre nœud.

Si un million de valeurs sont déjà réparties sur 100 nœuds, chaque nœud en additionne 10 000 localement. Une réduction en arbre combine les sommes partielles en sept tours, car $\lceil {\log }_{2}100\rceil =7$. Ces tours impliquent des messages réseau, pas seulement des additions locales. Garder les données près du travail qui les utilise réduit les transferts.

Les grappes de calcul passent à l’échelle avec des composants standards, et le matériel défaillant peut être remplacé. Une application peut toutefois nécessiter une reprise ou un redémarrage après la panne d’un nœud ; des nœuds remplaçables ne préservent pas à eux seuls la progression d’une tâche interrompue.

Mesurer plus que le débit des liaisons

La latence est le temps nécessaire à l’arrivée d’un message. La bande passante d’une liaison est son débit de données. La bande passante totale du réseau additionne la capacité des liaisons et décrit un total idéal où toutes les liaisons sont occupées. La bande passante de bissection additionne la capacité traversant la coupe de plus faible capacité qui divise les nœuds en deux moitiés égales.

Pour des liaisons de même capacité $B$, comparons des réseaux de 64 nœuds avec une même convention de comptage des liaisons :

Le maillage comporte huit lignes de sept liaisons horizontales et huit colonnes de sept liaisons verticales, soit 112. Relier les bords opposés dans les deux dimensions produit les 128 liaisons du tore. Chaque nœud de l’hypercube a six voisins, soit $64(6)\mathrm{/}2=192$ liaisons, car chaque liaison est comptée à ses deux extrémités.

Suivre la coupe, puis considérer le coût

Une coupe qui divise un anneau en deux moitiés ne traverse que deux liaisons. De nombreux messages peuvent donc se disputer cette frontière malgré les 64 liaisons de l’anneau au total. Dans un réseau totalement connecté, $32(32)=1024$ liaisons traversent la séparation entre les deux moitiés, mais le réseau nécessite $64(63)\mathrm{/}2=2016$ liaisons au total.

Davantage de liaisons peuvent améliorer la bande passante et réduire les sauts tout en augmentant le nombre de ports de commutation et de câbles ainsi que la puissance consommée. Les bus partagés disposent d’une seule ressource de communication partagée ; les maillages et les tores répartissent les liaisons dans l’espace. Les grandes grappes de calcul utilisent souvent des réseaux commutés à plusieurs étages, comme les arbres gras. Leur utilité dépend de la nature du trafic autant que de leur capacité totale. Pour la communication entre les moitiés d’une machine, la bissection révèle une limite que le total masque.

Un ordinateur à l’échelle d’un entrepôt associe des dizaines de milliers de serveurs à une alimentation électrique, un système de refroidissement et un réseau pour fournir des services internet. Les requêtes indépendantes fournissent du travail parallèle, tandis que des infrastructures de traitement par lots comme MapReduce répartissent de grandes tâches de traitement de données entre les machines.

Adapter la conception au service

Les systèmes à l’échelle d’un entrepôt privilégient le coût par requête, l’efficacité énergétique et la disponibilité. Les logiciels réessaient les tâches, répliquent les données et contournent les machines en panne. Un supercalculateur privilégie plus souvent un programme scientifique unique fortement couplé et une interconnexion rapide.

Les serveurs sont reliés à des commutateurs en haut de baie, eux-mêmes reliés par des niveaux supérieurs de commutation. Les niveaux supplémentaires ajoutent de la latence aux communications. Les liaisons montantes sont souvent sursouscrites, ce qui signifie que leur capacité est inférieure à la capacité totale des liaisons des serveurs. Certains réseaux non bloquants de type Clos offrent une capacité supérieure suffisante pour éviter cette restriction.

Quarante serveurs reliés à 10 Gbit/s offrent 400 Gbit/s à un commutateur de baie. Quatre liaisons montantes de 25 Gbit/s fournissent 100 Gbit/s, soit un rapport de 4:1. Si tous les serveurs émettent ensemble hors de la baie et se partagent la capacité également, chacun dispose en moyenne de 2,5 Gbit/s. Un serveur peu chargé peut utiliser davantage de bande passante lorsque les autres sont inactifs.

Séparer le surcoût de l’installation du calcul utile

L’indicateur d’efficacité énergétique, ou PUE, est l’énergie totale de l’installation divisée par l’énergie des équipements informatiques sur la même période. Une charge informatique constante de dix MW et un surcoût de 1,5 MW pour le refroidissement et la distribution donnent $PUE=\mathrm{11,5}\mathrm{/}10=\mathrm{1,15}$. La valeur idéale est de 1,0. Google a déclaré un PUE annuel de 1,09 pour l’ensemble de son parc en 2025.

Le PUE mesure le surcoût de l’installation, pas les calculs utiles par joule. Des serveurs inactifs ou inefficaces peuvent gaspiller de l’énergie parmi les équipements informatiques même lorsque le PUE de l’installation est excellent.

Pour 50 000 serveurs consommant en moyenne 300 W, la puissance informatique est de 15 MW. Avec un PUE de 1,2, la puissance de l’installation est de 18 MW. Maintenue pendant 8 760 heures, elle consomme 157 680 MWh ; le surcoût seul consomme 26 280 MWh. Les équipements électriques, le refroidissement, l’électricité, les serveurs et les bâtiments contribuent tous au coût de possession.

Une faible charge peut encore consommer beaucoup de puissance

La proportionnalité énergétique signifie que la puissance devrait augmenter proportionnellement à la charge. Un serveur qui consomme 150 W au repos et 300 W en crête, avec une augmentation linéaire, absorbe $150+\mathrm{0,3}(150)=195$ W à 30 % de charge. Il utilise 65 % de la puissance de crête pour 30 % du travail maximal. Des mesures antérieures à l’échelle d’entrepôts ont montré que de nombreux serveurs fonctionnaient entre environ 10 % et 50 % d’utilisation, ce qui rend la puissance au repos importante.

Prévoir les pannes et les réponses lentes

Si chacun des 50 000 serveurs connaît en moyenne une panne tous les dix ans, on attend 5 000 pannes par an dans le parc, soit environ 14 par jour. Cette prévision justifie la redondance et les nouvelles tentatives ; elle ne prévoit pas exactement 14 pannes chaque jour.

Une requête qui nécessite les réponses de nombreux serveurs attend la plus lente des réponses requises. Si chacun des 100 serveurs a indépendamment une probabilité de 1 % d’être lent, la probabilité qu’aucun ne le soit est ${\mathrm{0,99}}^{100}$. La probabilité qu’au moins l’un soit lent est donc $1-{\mathrm{0,99}}^{100}\approx \mathrm{0,634}$. La latence de queue d’un serveur devient une attente fréquente pour les requêtes lorsque leur diffusion touche de nombreux serveurs.