Parallele Prozessoren und Mehrkernprozessoren

Parallele Rechner führen Operationen gleichzeitig aus. Sie können die Latenz, also die Dauer einer Aufgabe, senken oder den Durchsatz, die Anzahl der pro Sekunde abgeschlossenen Aufgaben, erhöhen. Diese Ziele führen zu unterschiedlichen Arten der Arbeitsaufteilung.

Die Power Wall

Früher machte eine höhere Taktfrequenz jede Prozessorgeneration schneller. Die dynamische Leistungsaufnahme folgt näherungsweise $P\approx C{V}^{2}f$: Kapazität mal Spannung zum Quadrat mal Schaltfrequenz. Die Spannung geht in diesem Produkt doppelt ein. Senkt man Spannung und Frequenz jeweils um 15 %, ergibt sich ${\mathrm{0,85}}^2\times \mathrm{0,85}=\mathrm{0,614125}$ der ursprünglichen dynamischen Leistungsaufnahme, also etwa 61 %.

Die Spannung ließ sich nicht weiter ausreichend senken, um höhere Taktfrequenzen auszugleichen. Leckströme und Wärme machten eine weitere Frequenzsteigerung zu aufwendig. Diese Power Wall führte die Branche Mitte der 2000er-Jahre zu Mehrkernchips. Ein Mehrkernchip vereint mehrere Prozessorkerne auf einem Chip; ein Multiprozessorsystem besitzt mindestens zwei Prozessoren. Einfachere Kerne können pro Joule insgesamt mehr Arbeit leisten, sofern genügend Arbeit zur Aufteilung vorhanden ist.

Vier Arten paralleler Arbeit

ILP wird hauptsächlich von Hardware und Compilern erkannt. Pipelining, Mehrfachausgabe und Ausführung außerhalb der Programmreihenfolge überlappen Befehle und bewahren dabei die vorgeschriebenen Ergebnisse des Programms. Abhängigkeiten und Sprünge begrenzen die mögliche Überlappung.

DLP tritt auf, wenn viele Elemente dieselbe Berechnung benötigen. SIMD-Befehle, Vektorprozessoren und GPUs nutzen diese Regelmäßigkeit. Für TLP muss die Arbeit durch Code, eine Bibliothek oder einen Compiler auf Threads aufgeteilt werden. Ein Thread besitzt einen eigenen Befehlszähler, eigene Register und einen eigenen Stack; Threads eines Prozesses teilen sich dessen Adressraum. Getrennte Prozesse haben getrennte Adressräume.

Arbeit auf Anfrageebene benötigt wenig Zusammenarbeit zwischen den Aufgaben. Zusätzliche Server können die Kapazität daher nahezu proportional zu den hinzugefügten Ressourcen erhöhen. Auch mehrere unabhängige Programme auf einem Laptop nutzen sonst untätige Kerne und erhöhen den Gesamtdurchsatz, selbst wenn keines dieser Programme intern parallelisiert ist.

Nebenläufige Aufgaben und gleichzeitige Ausführung

Nebenläufigkeit bedeutet, dass Aufgaben während sich überschneidender Zeiträume Fortschritte machen. Ein Kern kann dies erreichen, indem er zwischen ihnen wechselt. Parallelität bedeutet, dass Operationen auf getrennter Hardware zum selben Zeitpunkt ausgeführt werden. Ein Einkernsystem kann mehrere Programme gleichzeitig in Bearbeitung halten, ohne ihre Befehle gleichzeitig auszuführen.

Ein Programm mit einem sequenziellen Thread verteilt sich nicht allein deshalb auf acht Kerne, weil diese vorhanden sind. Die zusätzlichen Kerne können andere Arbeit ausführen; um dieses Programm zu beschleunigen, muss es jedoch parallele Arbeit enthalten. Auch nach der Aufteilung begrenzt ein verbleibender sequenzieller Anteil den Gewinn.

Ein Prozessor kann viele Recheneinheiten enthalten, ohne vielen unabhängigen Befehlsströmen zu folgen. Die flynnsche Taxonomie klassifiziert Hardware anhand ihrer Befehls- und Datenströme. Daher wird zunächst festgelegt, welche Ebene des Rechners beschrieben wird.

Die vier Hardwareklassen

Ein herkömmlicher Einkernprozessor ist SISD, auch wenn seine Pipeline mehrere Befehle überlappt. Superskalare Befehlsausgabe erzeugt für sich genommen weder unabhängige Befehlszähler noch macht sie den Kern zu MIMD-Hardware.

Ein SIMD-Befehl könnte acht Paare von 32-Bit-Gleitkommazahlen in einem 256-Bit-Register addieren. Er besitzt acht Lanes für die Berechnung, die aber einem Befehl folgen. Zu den SIMD-Erweiterungen mit fester Breite gehören x86 SSE und AVX sowie Arm Neon; auch Vektorarchitekturen fallen in die SIMD-Klasse. Regelmäßige Arrayberechnungen eignen sich gut, weil viele Elemente dieselbe Operation benötigen.

MIMD-Prozessoren wählen ihre Befehle unabhängig voneinander. Mehrkernchips, Multiprozessoren mit gemeinsamem Speicher und Cluster gehören zu dieser Klasse. Die Kerne können unterschiedliche Programme ausführen oder Kopien desselben Programms, die zu unterschiedlichen Zeiten verschiedene Verzweigungen erreichen.

MISD vervollständigt die 2 × 2-Klassifikation. Es gibt wenige eindeutige praktische Beispiele. Manche fehlertoleranten Systeme und Verarbeitungspipelines werden so beschrieben, ihre Klassifikation hängt aber davon ab, wie ihre Ströme definiert werden.

Ein Programm kann dennoch viele Befehlsströme haben

SPMD, ein Programm für mehrere Datenbereiche, ist ein Programmierstil. Jede Arbeitseinheit führt eine Kopie desselben Programms aus und wählt anhand ihrer ID ihren Datenanteil. Die Arbeitseinheiten besitzen unabhängige Ausführungszustände, sodass SPMD üblicherweise auf MIMD-Hardware läuft.

Wenn vier Kerne 4.000 Werte summieren, bearbeitet der Kern mit der ID $i$ die Indizes $1000i$ bis $1000i+999$. Kern 0 bearbeitet 0 bis 999; Kern 3 bearbeitet 3000 bis 3999. Jeder erzeugt eine Teilsumme, und diese vier Summen werden zusammengeführt. Der gemeinsame Programmtext verlangt nicht, dass alle vier Kerne im selben Taktzyklus denselben Befehl ausgeben.

Jede Ebene gesondert klassifizieren

Betrachte eine CPU mit acht Kernen, deren Threads AVX-Befehle nutzen. Über die Kerne hinweg ist sie MIMD: Die Kerne schreiten unabhängig voneinander voran. Innerhalb einer AVX-Operation ist sie SIMD: Ein Befehl steuert mehrere Lanes. Würde man alle Recheneinheiten zusammenzählen, bliebe diese Struktur verborgen.

Ebenso ist ein Cluster, auf dem Kopien eines wissenschaftlichen Programms laufen, MIMD-Hardware mit SPMD-Software. Seine Knoten kommunizieren über Nachrichten. Eine GPU kombiniert viele unabhängig eingeplante Threadgruppen mit einer SIMD-ähnlichen Ausführung innerhalb jeder Gruppe; NVIDIA nennt diese threadbasierte Ausführungsart SIMT.

Um eine Beschreibung zu klassifizieren, frag dich, wer den nächsten Befehl auswählt. Frag dann, ob dieser Befehl ein Datenelement oder mehrere verarbeitet. Die Anzahl der Quelldateien, installierten Programme oder Recheneinheiten allein kann diese beiden Fragen nicht beantworten.

Ein Cache-Fehlzugriff kann einen Kern Hunderte Taktzyklen warten lassen. Hardware-Multithreading hält die Ausführungszustände mehrerer Threads in einem Kern vor, sodass während dieser Wartezeit nutzbare Befehle eines anderen Threads ausgeführt werden können.

Was wird vervielfacht und was gemeinsam genutzt?

Jeder Hardwarethread besitzt einen eigenen Befehlszähler und Registerzustand. Die Threads teilen sich die Ausführungseinheiten und Caches des Kerns sowie den Zugang zur Speicherbandbreite. Da die Zustände in der Hardware vorliegen, ist ein Wechsel viel schneller als das Sichern und Wiederherstellen eines Betriebssystemprozesses, was Tausende Taktzyklen kosten kann. Die Hardware kann in einem Taktzyklus oder weniger wechseln.

Das Betriebssystem sieht jeden Hardwarethread als logischen Prozessor. Acht physische Kerne mit zwei Hardwarethreads pro Kern stellen 16 logische Prozessoren bereit. Diese Zahl beschreibt verfügbare Threadkontexte und keine 16 vollständigen physischen Kerne.

Drei Planungsverfahren

Die feingranulare Planung kann reihum zwischen bereiten Threads wechseln, selbst wenn der vorherige Thread fortfahren könnte. Die grobgranulare Planung führt diesen Thread weiter aus, bis eine kostspielige Wartezeit einen Wechsel rechtfertigt. SMT nutzt einen Kern mit Mehrfachausgabe und dynamischer Befehlsplanung: Befehle verschiedener Threads können in einem Taktzyklus unterschiedliche freie Ausgabeplätze belegen. Registerumbenennung hält ihre Zustände getrennt. Intels Hyper-Threading ist eine SMT-Implementierung; einige IBM-POWER-Prozessoren unterstützen acht Hardwarethreads pro Kern.

Latenzverdeckung durchrechnen

Angenommen, ein Thread rechnet 20 Taktzyklen lang und wartet dann 200 Taktzyklen auf den Speicher. Sein Zyklus dauert 220 Taktzyklen. Allein belegt seine nutzbare Arbeit $20\mathrm{/}220\approx \mathrm{0,091}$, also 9,1 % der Zeit.

Mit vier Threads, frei versetzter Arbeit und ohne Wechselkosten oder zusätzliche Konkurrenz kann nutzbare Arbeit $4\times 20\mathrm{/}220\approx \mathrm{0,364}$, also 36,4 % der Zeit belegen. Der Kern hat weiterhin Lücken, weil vier Rechenphasen von jeweils 20 Taktzyklen die Wartezeiten nicht abdecken können.

Elf Threads liefern insgesamt 220 Taktzyklen Rechenarbeit. Während einer 200 Taktzyklen wartet, können die anderen zehn jeweils 20 Taktzyklen rechnen. Dieser ideale Ablauf hält den Kern ausgelastet. Die tatsächliche Rechenleistung hängt auch davon ab, ob die Threads um Speicher und Ausführungsressourcen konkurrieren.

Durchsatz und Abschlusszeit eines Threads

Multithreading erhöht vor allem den Durchsatz. Ein bestimmter Thread kann länger brauchen, weil andere Threads Taktzyklen und Cacheplatz nutzen, die er sonst erhalten würde. Zwei Hardwarethreads, die sich eine Recheneinheit teilen, können nicht beide die Kapazität eines zweiten vollständigen Kerns hinzugewinnen.

Wenn beide Threads bereits dieselben vollständig ausgelasteten Ausführungseinheiten beanspruchen, hat SMT wenig freie Kapazität zu füllen. Wartet einer, während der andere ausführen kann, ist die gemeinsame Nutzung hilfreich. GPUs verdecken Latenz in größerem Maßstab, indem sie für ihre Scheduler viele Threadgruppen bereithalten.

Ein paralleler Rechner ist nützlich, wenn er die Gesamtlaufzeit verändert. Der Speedup vergleicht den besten sequenziellen Lauf mit dem parallelen Lauf; die Effizienz setzt diesen Gewinn ins Verhältnis zur Anzahl der verwendeten Prozessoren.

Mit der gemessenen Zeit beginnen

Für $p$ Prozessoren gilt $S(p)=T_1\mathrm{/}{T}_p$ und $E(p)=S(p)\mathrm{/}p$. Eine sequenzielle Aufgabe von 100 s, die auf acht Kernen in 16 s abgeschlossen wird, hat den Speedup $100\mathrm{/}16=\mathrm{6,25}$. Ihre Effizienz beträgt $\mathrm{6,25}\mathrm{/}8=\mathrm{0,78125}$, also etwa 78 %. Miss für beide Läufe die Gesamtlaufzeit und verwende dieselbe feste Aufgabe.

Zeit in verbesserten und unveränderten Anteil aufteilen

Im amdahlschen Gesetz ist $f$ der Anteil der ursprünglichen Ausführungszeit, den eine Verbesserung beschleunigt. Es ist ein Zeitanteil und kein Anteil der Quellcodezeilen. Wird dieser Teil $s$-mal so schnell, gilt

Der unveränderte Teil benötigt weiterhin seine volle ursprüngliche Zeit. Für eine ideale parallele Aufteilung auf $p$ Prozessoren setze $s=p$. Dieses Modell schließt den Aufwand für Kommunikation, Lastungleichgewicht und Synchronisation aus; diese Kosten können die gemessene Zeit verlängern.

Bei 90 % paralleler Zeit und zehn Prozessoren beträgt die normierte Zeit $\mathrm{0,1}+\mathrm{0,9}\mathrm{/}10=\mathrm{0,19}$. Der Speedup ist $1\mathrm{/}\mathrm{0,19}\approx \mathrm{5,26}$ und die Effizienz etwa 53 %. Bei 100 Prozessoren beträgt die Zeit $\mathrm{0,1}+\mathrm{0,009}=\mathrm{0,109}$ und der Speedup etwa 9,17. Eine Verzehnfachung der Prozessorzahl hat den Speedup nicht verdoppelt.

Die Grenze vor dem Kauf weiterer Prozessoren bestimmen

Mit wachsender Prozessorzahl nähert sich die parallele Zeit der Null. Der sequenzielle Anteil bleibt bestehen und ergibt die Obergrenze $S_{max}=1\mathrm{/}(1-f)$. Ein sequenzieller Anteil von 10 % begrenzt den idealen Speedup auf 10; ein Anteil von 5 % begrenzt ihn auf 20. Endliche Prozessorzahlen nähern sich diesen Obergrenzen von unten.

Das amdahlsche Gesetz bewertet auch eine einzelne schnellere Befehlsklasse. Wenn die Multiplikation 30 % eines Laufs beansprucht und viermal so schnell wird, beträgt die normierte Zeit $\mathrm{0,7}+\mathrm{0,3}\mathrm{/}4=\mathrm{0,775}$. Das gesamte Programm ist nur $1\mathrm{/}\mathrm{0,775}\approx \mathrm{1,29}$-mal so schnell.

Von einem Zielwert rückwärts rechnen

Für einen Speedup von 90 auf 100 Prozessoren schreibe $1\mathrm{/}90=(1-f)+f\mathrm{/}100=1-\mathrm{0,99}f$. Daher gilt $f=(1-1\mathrm{/}90)\mathrm{/}\mathrm{0,99}\approx \mathrm{0,99888}$. Nur etwa 0,11 % der ursprünglichen Zeit dürfen sequenziell bleiben. Ein Ziel nahe der Prozessorzahl verlangt einen sehr kleinen unveränderten Anteil.

Ein gemessener Speedup kann manchmal die Prozessorzahl übersteigen. Durch Aufteilung einer großen Arbeitsmenge passt möglicherweise jeder Teil in den Cache, sodass die Kosten pro Operation sinken. Dieser superlineare Speedup verändert eine Voraussetzung des einfachen Modells: Die parallele Arbeit ist nicht mehr bloß die ursprüngliche Arbeit, die in gleiche Anteile geteilt wird. Prüfe Algorithmus, Datengröße und Cacheverhalten, bevor du eine solche Messung deutest.

Zusätzliche Prozessoren können entweder dieselbe Aufgabe verkürzen oder einen Rechner eine größere Aufgabe bearbeiten lassen. Starke Skalierung und schwache Skalierung messen diese unterschiedlichen Einsatzarten. Beide hängen davon ab, wie viel nutzbare Arbeit nach Abzug der Koordinationskosten übrig bleibt.

Feste und wachsende Arbeit

Bei starker Skalierung bleibt die Problemgröße fest. Bei schwacher Skalierung wächst die Arbeit proportional zur Prozessorzahl, sodass die Arbeit pro Prozessor konstant bleibt. Perfekte schwache Skalierung hält die Gesamtlaufzeit konstant, während die Gesamtarbeit wächst.

Betrachte zehn sequenzielle Additionen, gefolgt von der elementweisen Summe zweier Matrizen. Eine 10 × 10-Matrix benötigt 100 parallele Additionen. Die Gesamtzeit beträgt 110 Einheiten auf einem Prozessor, 20 auf zehn und 12 auf fünfzig: Die Speedups betragen 5,5 und etwa 9,2. Eine 40 × 40-Matrix benötigt 1.600 parallele Additionen. Die Zeiten betragen 1.610, 170 und 42 Einheiten: Die Speedups betragen etwa 9,5 und 38,3. Beim größeren Problem fällt die nutzbare parallele Arbeit gegenüber denselben zehn sequenziellen Additionen stärker ins Gewicht.

Gustafsons anderer Ausgangswert

Das gustafsonsche Gesetz nutzt den sequenziellen Zeitanteil, der beim parallelen Lauf gemessen wird. Bezeichnet man diesen Anteil mit $s$, beträgt der skalierte Speedup $S_G=s+(1-s)p$. Für $s=\mathrm{0,05}$ und 100 Prozessoren gilt $S_G=\mathrm{0,05}+\mathrm{0,95}(100)=\mathrm{95,05}$.

Das amdahlsche Gesetz ergibt bei 5 % sequenzieller Zeit im ursprünglichen Lauf auf einem Prozessor nur etwa 16,8 auf 100 Prozessoren. Diese Anteile haben unterschiedliche Bezugsgrößen und beantworten unterschiedliche Fragen. Ein skalierter Speedup von 95,05 bedeutet nicht, dass die ursprüngliche Aufgabe mit fester Größe 95,05-mal so schnell abgeschlossen wurde.

Die langsamste Arbeitseinheit bestimmt den Endzeitpunkt

Werden 400 gleich große Arbeitseinheiten auf acht Kerne verteilt, beträgt der ausgeglichene Anteil je 50. Angenommen, einer erhält 120, während die anderen sieben je 40 erhalten. Der Schritt dauert 120 Zeiteinheiten, sodass der Speedup $400\mathrm{/}120\approx \mathrm{3,33}$ beträgt. Wer durch die durchschnittliche Last dividiert, übersieht das Warten auf den am stärksten belasteten Kern.

Auch das Erzeugen und Einplanen von Threads, Nachrichten, Barrieren und Konkurrenz um gemeinsame Bandbreite verbrauchen Zeit. Granularität ist der Umfang der Berechnungen zwischen Kommunikations- oder Synchronisationsereignissen. Größere Teilstücke senken die Koordinationskosten pro nützlicher Operation, aber bei zu wenigen Teilstücken können einige Arbeitseinheiten untätig bleiben.

Mehr Prozessoren können die Zeit verlängern

Für eine vollständig parallelisierbare Aufgabe von 100 s mit einem Aufwand von $\mathrm{0,5}p$ Sekunden gilt $T_p=100\mathrm{/}p+\mathrm{0,5}p$. Zehn Prozessoren benötigen 15 s. Vierzehn benötigen etwa 14,14 s, nahe dem Minimum bei ungefähr $\sqrt{200}$ Prozessoren. Vierzig benötigen 22,5 s, weil ihr Aufwand von 20 s die kürzere nutzbare Berechnung überwiegt.

Gib bei der Ergebnisdarstellung an, ob die Arbeit fest blieb oder wuchs, und vergleiche mit der Zeit des besten sequenziellen Programms. Wenn du eine langsame parallele Implementierung auf einem Prozessor als Ausgangswert nutzt, kann der angegebene Speedup überhöht sein.

Die beworbene Rechenrate eines Prozessors ist nur eine Grenze. Das Roofline-Modell fragt auch, ob der Speicher Daten schnell genug liefern kann, um diese Rate aufrechtzuerhalten. Es gibt eine obere Schranke an und hilft damit, eine sinnvolle Richtung für Optimierungen zu finden.

Operationen pro DRAM-Byte zählen

Die arithmetische Intensität ist die Anzahl der Gleitkommaoperationen geteilt durch die zwischen DRAM und Prozessor übertragenen Byte. Die ursprüngliche Arbeit zum Roofline-Modell nennt dieses DRAM-bezogene Maß operationale Intensität. Zähle den Datenverkehr an dieser Schnittstelle einschließlich der Lese- und Schreibzugriffe.

DAXPY berechnet $y_i=a{x}_i+y_i$ mit Zahlen doppelter Genauigkeit. Jedes Element benötigt eine Multiplikation und eine Addition: zwei Gleitkommaoperationen. Das Lesen eines 8 Byte großen Werts aus jedem Array und das Zurückschreiben von 8 Byte übertragen 24 Byte. Die Intensität beträgt $I=2\mathrm{/}24\approx \mathrm{0,0833}$ Gleitkommaoperationen pro Byte. Der skalare Koeffizient wird wiederverwendet; in diesem Verkehrsmodell wird er nicht für jedes Element erneut aus dem DRAM geladen.

Beide Grenzen anwenden

Sind die Spitzenrechenleistung $P$ und die maximale Speicherbandbreite $B$, dann ist das Dach $\min (P,BI)$. Bandbreite mal Gleitkommaoperationen pro Byte ergibt die Einheit Gleitkommaoperationen pro Sekunde. In einem doppelt logarithmischen Diagramm steigt die Bandbreitengrenze an und die Rechenleistungsgrenze verläuft waagerecht.

Bei 1.000 GFLOP/s und 100 GB/s liegt der Knickpunkt bei $P\mathrm{/}B=1000\mathrm{/}100=10$ Gleitkommaoperationen pro Byte. Unterhalb dieser Intensität liegt das Bandbreitendach niedriger, darüber das Rechenleistungsdach. DAXPY ist durch $100(2\mathrm{/}24)\approx \mathrm{8,33}$ GFLOP/s begrenzt, also weniger als 1 % der Spitzenrechenleistung des Rechners. Ein Kernel mit der Intensität 2 ist auf 200 GFLOP/s begrenzt.

Ein speichergebundener Kernel kann dennoch von zusätzlichen Kernen oder breiteren SIMD-Einheiten profitieren, solange zu wenige Anfragen gleichzeitig bearbeitet werden, um die verfügbare Bandbreite zu nutzen. Sobald die Bandbreite ausgeschöpft ist, hebt zusätzliche Rechenkapazität sein Dach nicht an. Eine höhere Bandbreite oder mehr nutzbare Berechnung pro übertragenem Byte hebt es dagegen an.

Wiederverwendung verändert die Intensität

Die Multiplikation zweier 1000 × 1000-Matrizen mit doppelter Genauigkeit führt etwa $2(1000{)}^3=2\times {10}^9$ Gleitkommaoperationen aus. Wenn jedes Element der beiden Eingabematrizen und der Ausgabematrix die DRAM-Grenze einmal überquert, beträgt der gesamte Datenverkehr 24 MB. Das ergibt etwa 83 Gleitkommaoperationen pro Byte und liegt über dem Knickpunkt des Beispielrechners.

Eine naive Implementierung kann Matrixwerte wiederholt aus dem DRAM lesen. Kachelung organisiert Schleifen um Teilmatrizen, die in den Cache passen, sodass Werte für viele Operationen genutzt werden, bevor sie ersetzt werden. Der Algorithmus bleibt eine Matrixmultiplikation, aber die Implementierung überträgt pro Operation weniger Byte aus dem oder in das DRAM.

Ein Dach ist kein Messergebnis

Ein Kernel kann wegen Abhängigkeiten, Befehlsaufwand, ungünstiger Zugriffsmuster oder zu wenig paralleler Arbeit unter der Schranke bleiben. Niedrigere Leistungsgrenzen zeigen solche Einschränkungen: Fehlende SIMD-Befehle können die Rechenleistungsgrenze senken, ungünstige Zugriffsmuster dagegen die nutzbare Bandbreite. Ein Punkt unterhalb eines Dachs muss untersucht werden; das Dach allein sagt seine genaue Rechenrate nicht voraus.

Eine SIMD-Erweiterung teilt ein breites Register in Lanes. Ein Befehl führt in jeder Lane dieselbe Operation aus. Deshalb eignen sich regelmäßige Schleifen über unabhängige Elemente gut dafür.

Lanes und Iterationen zählen

Die Anzahl der Lanes ist die Registerbreite geteilt durch die Elementbreite. Ein 256-Bit-Register fasst acht 32-Bit-Gleitkommazahlen einfacher Genauigkeit, vier 64-Bit-Gleitkommazahlen doppelter Genauigkeit oder 32 Elemente zu je einem Byte. SSE verwendet 128-Bit-XMM-Register; AVX und AVX2 verwenden 256-Bit-YMM-Register. Im 64-Bit-Modus gibt es 16 YMM-Register. AVX-512 stellt 32 512-Bit-ZMM-Register bereit, die jeweils 16 Gleitkommazahlen einfacher Genauigkeit fassen. Arm Neon verwendet 128-Bit-Register.

Für c[i] = a[i] + b[i] über 1.000 Gleitkommazahlen verarbeiten Befehle mit acht Lanes das Array in 125 Vektoriterationen. Jede Iteration lädt zwei Vektoren, addiert deren Lanes und speichert einen Vektor. Bei 1.003 Elementen bleiben nach 125 vollständigen Iterationen drei Elemente übrig. Eine skalare Restschleife oder eine maskierte Operation verarbeitet diesen kurzen Rest, ohne über das Ende der Arrays hinaus zuzugreifen.

Acht Lanes beschreiben eine Rechenkapazität und garantieren keinen achtfachen Speedup des Programms. Initialisierung, Restverarbeitung, sequenzielle Arbeit und Speicherbandbreite beanspruchen weiterhin Zeit.

Eine Anordnung wählen, die ein Register füllt

Einfache Vektorladebefehle eignen sich gut für zusammenhängende Elemente mit Schrittweite 1. Ein Array aus {x, y, z}-Punkten mit 4 Byte pro Koordinate legt aufeinanderfolgende x-Felder im Abstand von 12 Byte ab. Eine Struktur von Arrays legt die x-Werte zusammen ab, danach die y-Werte und dann die z-Werte. Werden nur die x-Werte verarbeitet, kann nun ein zusammenhängender Ladezugriff einen Vektor füllen. Gather- und Scatter-Befehle unterstützen verstreute Adressen, verursachen aber meist höhere Zugriffskosten.

Bedingte Arbeit innerhalb der Lanes ausführen

Für c[i] = (a[i] > 0) ? a[i] : 0 erzeugt ein Vergleich eine Maske mit einem Wert pro Lane. Ein Auswahlbefehl wählt zwischen möglichen Ergebnissen, oder maskierte Befehle aktualisieren nur die ausgewählten Lanes. Inaktive Lanes behalten oder erhalten Werte entsprechend dem Maskenverhalten des Befehls.

Wenn unterschiedliche Lanes lange, unterschiedliche Verzweigungen benötigen, müssen die Befehle dieser Pfade weiterhin ausgegeben werden, wobei nur die jeweils betroffenen Lanes aktiv sind. Die Gruppe kann nicht jeder Lane einen eigenen Befehlsstrom geben. Eine Maske erhält die Korrektheit, lässt aber nicht alle Befehle dieser Pfade verschwinden.

Unabhängigkeit der Iterationen prüfen

Compiler können Schleifen automatisch vektorisieren; Programmierer können auch Bibliotheken oder Intrinsics wie _mm256_add_ps verwenden. Der Compiler muss nachweisen, dass gleichzeitige Iterationen das Ergebnis bewahren. Die Rekurrenz a[i] = a[i-1] + b[i] hängt vom neuen Wert der vorhergehenden Iteration ab und kann nicht direkt als unabhängige Elemente vektorisiert werden. Eine mögliche Überlappung von Zeigern, Aliasing genannt, kann ebenfalls eine Überlappungsprüfung zur Laufzeit nötig machen oder die Vektorisierung verhindern.

Eine Reduktion hält mehrere Teilsummen in einem Register und führt anschließend die Lanes zusammen. Das Umordnen von Gleitkommaadditionen kann die letzten Bits verändern, weil die Addition nicht assoziativ ist. Ganzzahlige Summen stimmen überein, wenn bei der gewählten Arithmetik kein Überlauf auftritt oder definierte modulare Arithmetik verwendet wird.

SIMD und Mehrkernausführung lassen sich kombinieren. Acht Kerne, die jeweils eine Addition mit acht Lanes pro Taktzyklus ausgeben, können bis zu 64 Gleitkommaadditionen pro Taktzyklus ausführen. Sie unterliegen weiterhin gemeinsamen Grenzen wie der Speicherbandbreite.

Eine Vektorarchitektur beschreibt eine Operation auf einer Folge von Elementen und speichert die aktive Vektorlänge in einem Register. Dasselbe Programm kann auf Implementierungen mit unterschiedlich großen Vektorregistern laufen.

Kapazität, Länge und Lanes unterscheiden

Die RISC-V-Vektorerweiterung RVV wurde 2021 in Version 1.0 ratifiziert. Sie stellt 32 Vektorregister bereit, von v0 bis v31. Jede Implementierung wählt ihre Registerbreite VLEN als Zweierpotenz. Elementbreite und Registergruppierung bestimmen, wie viele Elemente hineinpassen.

Die aktive Länge vl ist die Anzahl der Elemente, die die nächsten Vektorbefehle bearbeiten. vsetvli konfiguriert die Vektoroperation anhand der verbleibenden Elementzahl und des gewünschten Elementformats und gibt die gewählte Länge vl zurück. Der Code rückt um diesen Rückgabewert weiter. Das ist wichtig, weil portabler Code die tatsächlich von der Implementierung gewählte Länge verwenden muss.

Lanes sind die parallelen Pipelines, die die Vektoroperation ausführen. Eine Addition mit 32 Elementen auf vier Lanes benötigt nach dem Anlaufen der Pipeline etwa acht Taktzyklen. Ein Vektor mit 32 Elementen benötigt keine 32 Recheneinheiten. Längere Vektoren und mehr Lanes sind voneinander unabhängige Entwurfsentscheidungen.

Eine Schleife mit Strip-Mining aufteilen

Strip-Mining teilt eine lange Schleife in vektorgroße Stücke. Für 100 Elemente und höchstens 32 pro Operation ist 32, 32, 32 und 4 eine mögliche Folge von Längen. Jeder Durchlauf lädt den nächsten Teil, berechnet und speichert ihn und verringert die verbleibende Anzahl um den Wert von vl.

DAXPY über 1.000 Gleitkommazahlen doppelter Genauigkeit kann entsprechend 31 Durchläufe mit je 32 und einen mit acht Elementen verwenden, insgesamt also 32 Durchläufe. RVV kann gegen Ende andere zulässige Längen wählen. Deshalb verwendet die Schleife den zurückgegebenen Wert von vl, statt eine bestimmte letzte Aufteilung anzunehmen. Das Setzen der aktiven Länge behandelt den kurzen Rest ohne separate skalare Restschleife.

Das Muster zur Schleifensteuerung lautet:

1. Fordere eine Länge für die verbleibenden Elemente an.
2. Lade und berechne diese Anzahl von Elementen.
3. Speichere die Ergebnisse.
4. Setze die Zeiger weiter und ziehe die zurückgegebene Länge ab.

Ein Vektorbefehl beschreibt viele Elementoperationen. Gegenüber einem skalaren Rumpf, der für jedes Element wiederholt wird, verringert das den Aufwand für das Holen und Dekodieren von Befehlen.

Abhängige Operationen und nicht zusammenhängende Daten versorgen

Verkettung erlaubt einem verarbeitenden Vektorbefehl, mit frühen Ergebnissen zu beginnen, bevor der erzeugende Befehl spätere Elemente fertiggestellt hat. Bei einer Multiplikation mit anschließender Addition kann das erste Produkt in die Additionspipeline gelangen, während spätere Produkte noch gebildet werden.

Eine Maske, bei RVV häufig im Register v0, wählt aus, welche Elemente aktualisiert werden. Ein Ladebefehl mit fester Schrittweite greift mit einem festen Byteabstand zu. In einer zeilenweise gespeicherten 100 × 100-Matrix aus Gleitkommazahlen doppelter Genauigkeit führt ein Schritt nach unten in einer Spalte um 100 Elemente oder 800 Byte weiter. Indizierte Lade- und Speicherbefehle verwenden explizite Adressen für unregelmäßige Daten.

Sowohl Vektorarchitekturen als auch Erweiterungen fester Breite gehören in der flynnschen Taxonomie zu SIMD. Erweiterungen fester Breite kodieren bestimmte Breiten; RVV-Code beschreibt Längen, die sich an die verfügbare Hardware anpassen. Beide benötigen weiterhin eine Ausführungsreihenfolge, die die Datenabhängigkeiten der Schleife beachtet.

Eine CPU konzentriert ihre Ressourcen auf geringe Latenz für wenige Threads. Eine GPU konzentriert ihre Ressourcen auf hohen Durchsatz über viele datenparallele Threads. Bei der Wahl zwischen beiden muss die gesamte Aufgabe einschließlich der Datenübertragung betrachtet werden.

Threads laufen auf Streaming Multiprocessors

Eine GPU von NVIDIA enthält viele Streaming Multiprocessors, kurz SMs. Ein SM besitzt Rechen-Lanes, Warp-Scheduler, einen großen Registersatz und schnellen Shared Memory auf dem Chip. Ein „CUDA Core“ bezeichnet eine Rechen-Lane und keinen vollständigen unabhängigen Kern wie bei einer CPU. AMD nennt die entsprechende Verarbeitungsgruppe Compute Unit.

NVIDIA fasst Threads zu Warps aus jeweils 32 Threads zusammen. Ein ausgegebener Befehl wirkt auf die aktiven Threads des Warps, die jeweils eigene Registerwerte und Daten besitzen. Das ist SIMT: ein einzelner Befehl für mehrere Threads. AMD verwendet Wavefronts mit je nach Architektur 32 oder 64 Threads.

Ein SM hält viele Warps vor. Wartet einer auf den Speicher, wählt der Scheduler einen anderen bereiten Warp. Mehr bereite Threads als Lanes verdecken Latenz; sie verringern nicht die Latenz eines einzelnen DRAM-Zugriffs.

Divergenz lässt Lanes untätig

Nehmen Threads eines Warps unterschiedliche Verzweigungen, führt der Warp die benötigten Pfade aus und maskiert dabei die jeweils anderen Threads. Eine gleichmäßig aufgeteilte if-else-Verzweigung mit gleich aufwendigen Zweigen kann ungefähr die doppelte Zeit eines Pfads benötigen, weil beide Pfade ausgeführt werden müssen. Eine Verzweigung kann unproblematisch sein, wenn ihre Bedingung innerhalb jedes Warps einheitlich ist, selbst wenn verschiedene Warps unterschiedliche Pfade wählen.

Speicherzugriffe aufeinander abstimmen

Register pro Thread halten private Werte. Shared Memory pro Block ist ein vom Programm verwalteter Speicher auf dem Chip. L1- und L2-Caches ermöglichen automatische Wiederverwendung, während der globale Speicher das DRAM der GPU ist, etwa GDDR oder HBM.

Wenn 32 Threads aufeinanderfolgende Wörter zu je 4 Byte lesen, kann ihr Bereich von 128 Byte durch wenige breite Transaktionen bedient werden. Dieses Coalescing nutzt die Bandbreite effizient. Verstreute Adressen können für dieselben nutzbaren Daten viele Transaktionen benötigen. Die Wiederverwendung im Shared Memory kann wiederholten Verkehr zum globalen Speicher weiter verringern.

Die Übertragungszeit einbeziehen

Bei einer separaten GPU werden Eingaben und Ergebnisse über eine Verbindung wie PCIe oder NVLink übertragen. Das Senden von 800 MB und Empfangen von 400 MB überträgt bei effektiv 20 GB/s insgesamt 1,2 GB in 0,06 s, also 60 ms. Mit einem Kernel von 10 ms ergeben sich insgesamt 70 ms. Eine CPU, die in 50 ms fertig ist, gewinnt diesen Vergleich trotz der viel schnelleren GPU-Berechnung.

Bleiben Daten über mehrere Kernel hinweg auf der GPU, entfallen wiederholte Übertragungen. Große, regelmäßige Aufgaben können außerdem genügend nutzbare Rechenarbeit bieten, um die Kosten für Übertragung und Start aufzuwiegen. Bei kleinen, sequenziellen Aufgaben mit vielen Verzweigungen ist oft eine CPU im Vorteil.

Bei geeigneter Arbeit weiter spezialisieren

Ein heterogenes System verwendet die CPU für die Ausführung von Programm und Betriebssystem und einen Beschleuniger für geeignete Teile. Domänenspezifische Beschleuniger wie TPUs und NPUs widmen ihre Hardware einer enger gefassten Aufgabe, häufig der Matrixmultiplikation mit 8-Bit-Ganzzahlen oder 16-Bit-Gleitkommazahlen. Spezialisierung kann in diesem Bereich die Rechenleistung pro Watt verbessern und verringert zugleich die Bandbreite der Aufgaben, die die Hardware gut bewältigt.

Ein GPU-Kernel läuft einmal pro gestartetem Thread. Der Host startet ein Grid aus Threadblöcken, und jeder Thread berechnet, welchen Teil der Daten er bearbeitet. CUDA nennt die CPU den Host und die GPU das Gerät; OpenCL verwendet die verwandten Begriffe Work-Item und Work-Group.

Einen eindeutigen Index berechnen

Bei einem eindimensionalen Start lautet der Elementindex blockIdx.x * blockDim.x + threadIdx.x. Bei 256 Threads pro Block bearbeitet Thread 5 in Block 2 das Element 517: Die beiden vorhergehenden Blöcke umfassen 512 Elemente, und der Thread ergänzt einen Versatz von fünf. Thread- und Blockindizes beginnen bei null.

Ein SAXPY-Kernel wendet eine Multiplikation und eine Addition auf jede Gleitkommazahl an:

__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];
    }
}

Das Grid ersetzt die gewöhnliche Schleife über die Elemente. Jeder aktive Thread aktualisiert sein eigenes y[i]; benachbarte Threads verwenden benachbarte Adressen und unterstützen so Coalescing.

Die Blockzahl aufrunden

Für 1.000.000 Elemente und 256 Threads pro Block benötigt das Grid $\lceil 1000000\mathrm{/}256\rceil =3907$ Blöcke. Diese starten 1.000.192 Threads. Die letzten 192 haben Indizes außerhalb des Arrays, sodass die Bedingung i < n ihre Speicherzugriffe verhindert. Abrunden würde zu wenige Threads starten und gültige Elemente unbearbeitet lassen.

Ein Block mit 256 Threads bildet acht Warps. Blöcke können beliebigen SMs zugewiesen und in beliebiger Reihenfolge ausgeführt werden. Die Korrektheit des Programms darf nicht davon abhängen, dass Block 0 vor Block 1 fertig wird.

Arbeit von Host und Gerät organisieren

Ein typischer Host reserviert Gerätespeicher, kopiert die Eingaben, startet saxpy<<<3907, 256>>>(n, 2.0f, d_x, d_y) und erhält die Ergebnisse. Der Start ist asynchron: Die CPU kann fortfahren, während der Kernel läuft. Der Host muss durch Synchronisation oder eine geeignete blockierende Ergebniskopie warten, bevor er die fertigen Daten verwendet.

Innerhalb eines Blocks zusammenarbeiten

Threads eines Blocks können Werte über Shared Memory austauschen und sich bei __syncthreads() treffen. Jeder Thread, der an dieser blockweiten Barriere teilnimmt, muss sie erreichen. Die Barriere innerhalb von if (i < n) zu platzieren, ist unsicher, wenn nur einige Threads die Bedingung erfüllen.

Für eine Reduktion von 1.024 Werten im Shared Memory addieren 512 Threads Paare, dann führen 256 die Ergebnisse zusammen, danach 128 und so weiter. Zehn Halbierungsstufen lassen eine Summe übrig. Eine Barriere zwischen den Stufen verhindert, dass ein Thread ein Teilergebnis verarbeitet, bevor der erzeugende Thread es schreibt.

Ein gewöhnlicher Start besitzt keine gridweite Barriere. Getrennte Blöcke können Teilsummen schreiben, die ein zweiter Kernel zusammenführt; alternativ können geeignete atomare Operationen ein globales Ergebnis aktualisieren. Eine Blockbarriere allein kann verschiedene Blöcke nicht koordinieren. Eine Schleife mit Grid-Schrittweite ist ein weiteres Zuordnungsmuster: Jeder Thread bearbeitet mehrere Elemente und rückt dabei jeweils um die gesamte Grid-Größe weiter.

Ein Multiprozessor mit gemeinsamem Speicher stellt seinen Prozessoren einen physischen Adressraum zur Verfügung. Threads kommunizieren durch Lade- und Speicherzugriffe auf gemeinsame Variablen und halten zugleich private Werte wie lokale Teilsummen vor.

Gemeinsame Adressen können unterschiedlich viel kosten

Bei UMA, dem gleichförmigen Speicherzugriff, ist die Speicherzugriffszeit von jedem Prozessor aus ungefähr gleich. Bei NUMA, dem nicht gleichförmigen Speicherzugriff, erreicht ein Sockel seinen lokalen Speicher schneller als den Speicher eines anderen Sockels. Der Adressraum wird gemeinsam genutzt, obwohl der physische Speicher verteilt ist.

Bei 70 % lokalen Zugriffen mit 80 ns und 30 % entfernten Zugriffen mit 140 ns beträgt das gewichtete arithmetische Mittel $\mathrm{0,7}(80)+\mathrm{0,3}(140)=98$ ns. Werden Daten nahe bei ihrem Hauptnutzer abgelegt, sinkt der Anteil entfernter Zugriffe. Würde man jeden Zugriff als lokal zählen, würde man die Zeit unterschätzen.

Eine verbreitete Cacheanordnung gibt jedem Kern private L1-Befehls- und Datencaches, häufig einen privaten L2-Cache und einen gemeinsam genutzten Cache der letzten Ebene. Ein Ring oder Gitter verbindet die Kerne und Caches. Gemeinsamer Speicher beschreibt die Adressierung; er verlangt nicht, dass jeder Kern denselben Cache nutzt.

Aufteilen, berechnen und reduzieren

Um 64.000 Werte auf 64 Kernen zu summieren, bearbeitet jeder Kern 1.000 Werte. Die 64 Teilsummen können in einem Reduktionsbaum zusammengeführt werden. Die Anzahl aktiver Addierer halbiert sich auf jeder Stufe: 32, 16, 8, 4, 2 und 1. Sechs Stufen führen die Teilergebnisse zusammen.

Mit etwa 1.000 lokalen Additionsschritten und sechs Schritten zum Zusammenführen ergeben sich vor Berücksichtigung des Aufwands etwa 1.006 Schritte auf dem längsten Pfad. Der Baum vermeidet, dass jede Eingabe einen gemeinsamen Zähler aktualisiert. Jede Stufe hängt von den fertigen Werten der vorherigen Stufe ab, sodass sich die Arbeitseinheiten zwischen den Stufen an Barrieren treffen.

OpenMP kann dies mit #pragma omp parallel for reduction(+:sum) ausdrücken. Die Anweisung teilt Iterationen auf, gibt Threads private Teilsummen und führt sie zusammen. POSIX-Threads bieten eine explizitere Kontrolle über die Threadverwaltung.

Eine Aktualisierung unteilbar machen

Eine atomare Lese-Änderungs-Schreiboperation macht eine Aktualisierung aus Sicht anderer Prozessoren unteilbar. Atomarer Austausch und Fetch-and-Add sind Beispiele. Ein Laden mit anschließendem gewöhnlichem Speichern sind zwei getrennte Operationen, sodass ein anderer Thread dazwischen eingreifen kann.

RISC-V stellt lr.w und sc.w bereit. Der reservierende Ladebefehl liest ein Wort und legt eine Reservierung an. Der bedingte Speicherbefehl versucht zu schreiben und meldet Erfolg oder Misserfolg. Ein zwischenzeitlicher Schreibzugriff auf die reservierte Stelle lässt den Versuch fehlschlagen; die Software versucht es erneut. Die Reservierung blockiert keine anderen Prozessoren, und ein Fehlschlag kann auch aus anderen zulässigen Gründen auftreten. Diese Grundoperationen unterstützen Synchronisationsmechanismen.

Gemeinsame Engpässe beachten

Kann jeder Kern Daten mit 15 GB/s kontinuierlich verarbeiten, der Sockel aber nur 60 GB/s liefern, schöpfen etwa vier solche Datenströme die Bandbreite aus. Mehr Kerne können die Datenrate des Sockels nicht erhöhen. Auch Kommunikation über Caches verursacht Aufwand: Ein schreibender Kern ohne Schreibberechtigung muss den Besitz erlangen und andere Kopien invalidieren. Ein Kern, der bereits einen Block in Modified hält, kann lokal weiterschreiben; spätere Leser müssen möglicherweise eine neue Kopie holen.

Private Caches verringern die Latenz, können aber mehrere Kopien derselben Speicherstelle halten. Cachekohärenz verhindert, dass diese Kopien widersprüchliche Folgen von Schreibzugriffen auf diese Stelle zeigen.

Warum die Aktualisierung des Speichers nicht genügt

Zu Beginn steht X = 0 im Speicher. Kern A liest den Wert null und speichert ihn im Cache; Kern B tut dasselbe. A schreibt den Wert eins in seinen Cache mit Rückschreibverfahren. Ohne Kohärenzprotokoll kann B einen Treffer im eigenen Cache haben und weiterhin null lesen. Das Durchschreiben in den Speicher würde Bs bereits zwischengespeicherte Kopie nicht korrigieren.

Kohärenz arbeitet auf der Ebene von Cacheblöcken, üblicherweise 64 Byte. Ein Controller verfolgt den Zustand jedes Blocks und befolgt Protokollregeln, wenn Prozessoren ihn lesen oder schreiben. Migration verschiebt nutzbare Daten in den Cache ihres aktuellen Nutzers. Replikation ermöglicht mehreren Lesern, unveränderte Kopien zu halten. Kohärenz muss beide Vorteile bewahren und zugleich Schreibzugriffe steuern.

Drei erforderliche Eigenschaften

Wird auf eine Speicherstelle erst der Wert eins und danach der Wert zwei geschrieben, können Leser nicht unterschiedlicher Meinung darüber sein, welcher Schreibzugriff später kam. Erst den späteren und dann den früheren Schreibzugriff zu beobachten, würde diese Reihenfolge umkehren. Ein neuer nachfolgender Schreibzugriff könnte natürlich erneut den Wert eins speichern; Serialisierung betrifft Schreibereignisse und nicht die Frage, ob ein Zahlenwert jemals wieder auftreten darf.

Besitz und Kopien steuern

In einem Schreibinvalidierungsprotokoll erlangt ein schreibender Kern exklusive Schreibberechtigung und invalidiert andere Cachekopien. Ein späterer Lesezugriff auf einen invalidierten Block führt zu einem Fehlzugriff und holt aktuelle Daten. Dies ist ein kohärenzbedingter Fehlzugriff, zusätzlich zu obligatorischen, kapazitätsbedingten und konfliktbedingten Fehlzugriffen.

Ein Schreibaktualisierungsprotokoll sendet neue Werte stattdessen an vorhandene Kopien. Wiederholte Schreibzugriffe können dies aufwendig machen. Snooping sendet Kohärenzanfragen an alle, damit Caches reagieren können. Ein Kohärenzverzeichnis erfasst Mitnutzer und leitet Nachrichten an die betreffenden Caches.

Kohärenz, Reihenfolge und Atomarität

Kohärenz betrifft eine Speicherstelle. Ein Speicherkonsistenzmodell definiert zulässige beobachtbare Reihenfolgen und Ergebnisse von Lese- und Schreibzugriffen, einschließlich Operationen auf verschiedenen Stellen. Sequenzielle Konsistenz verlangt ein Ergebnis, das einer Verschachtelung entspricht, die die Programmreihenfolge jedes Prozessors bewahrt. Schwächere Modelle erlauben mehr Umordnungen; Fence-Befehle erzwingen bei Bedarf die erforderliche Reihenfolge.

Weder kohärente Kopien noch eine konsistente Reihenfolge machen eine aus mehreren Befehlen bestehende Erhöhung atomar. Angenommen, zwei Threads erhöhen einen gemeinsamen Zähler, der anfangs fünf beträgt. Eine mögliche Folge ist: A liest fünf, B liest fünf, A schreibt sechs, B schreibt sechs. Jeder Lese- und Schreibzugriff kann kohärent sein, dennoch beträgt der Endstand sechs statt sieben. Das Programm benötigt eine atomare Aktualisierung oder geeignete Synchronisation, um beide Erhöhungen zu erhalten.

Ein Snooping-Cachecontroller überwacht Kohärenzanfragen auf einer Verbindung mit Rundsendung. Um ein Protokoll zu verfolgen, halte für die Blockkopie jedes Kerns einen Zustand fest und aktualisiere nach jedem Zugriff alle betroffenen Kopien.

Zustände als Berechtigungen lesen

Im grundlegenden MSI sendet ein Lesefehlzugriff eine Busleseanfrage und lädt den Block in S. Das Schreiben eines Blocks in S benötigt ein Upgrade, um andere Kopien zu invalidieren. Das Schreiben aus I benötigt eine exklusive Leseanfrage, um Daten zu holen und Besitz zu erlangen. Der schreibende Kern endet in M. Ein Treffer in M benötigt keine Kohärenztransaktion.

Fünf Zugriffe verfolgen

Beginne mit X = 0 im Speicher und A und B jeweils in Invalid.

1. A liest X. Eine Busleseanfrage liefert null. Zustände: A = S, B = I.
2. B liest X. Eine weitere Busleseanfrage liefert null. Zustände: A = S, B = S.
3. A schreibt eins. Sein Upgrade invalidiert B. Zustände: A = M, B = I.
4. B liest X. B hat einen Fehlzugriff; A liefert eins und wechselt nach S. Im grundlegenden MSI wird auch der Speicher wieder aktuell. Zustände: A = S, B = S.
5. B schreibt zwei. Sein Upgrade invalidiert A. Zustände: A = I, B = M.

Jeder Zugriff in diesem Ablauf benötigt eine Bustransaktion. Der vierte Zugriff ist die wichtige Prüfung des Besitzes: Der Speicher allein enthielt noch den alten Wert, bevor A seine veränderten Daten lieferte.

Was Exclusive spart

MESI ergänzt E für eine einzelne unveränderte Cachekopie. Liest A einen nicht zwischengespeicherten Block und hat kein anderer Cache ihn, erhält A den Zustand E. Ein anschließender Schreibzugriff ändert E stillschweigend in M, weil keine anderen Kopien invalidiert werden müssen. MSI würde zunächst S laden und für diesen Schreibzugriff dann ein Upgrade senden.

Liest ein anderer Kern, während A den Zustand E hält, wechseln beide nach S. M und E haben beide einen gültigen Besitzer, aber nur M enthält Daten, die neuer als der Speicherinhalt sind.

Wiederholte Schreibzugriffe und Skalierung

Ein schreibender Kern, der M erlangt hat, kann den Block zehnmal aktualisieren, ohne zehn neue Invalidierungen auszulösen. Ein Schreibaktualisierungsverfahren würde die geänderten Werte an die Mitnutzer senden.

Snooping mit Rundsendung wird aufwendig, wenn immer mehr Kerne jede Anfrage überwachen. Verzeichnisstrukturen verringern diese Auffächerung. Ausgefeiltere Protokolle ergänzen nützliche Zustände: Der Owned-Zustand von MOESI erlaubt eine veränderte gemeinsam genutzte Kopie, und MESIF bestimmt einen Mitnutzer im Forward-Zustand zur Beantwortung von Anfragen.

Jede Kohärenzanfrage an alle zu senden, wird auf großen Rechnern aufwendig. Ein Kohärenzverzeichnis erfasst, welche Caches einen Block möglicherweise besitzen, sodass Anfragen an diese Caches gehen können. Dieselbe Erfassung auf Blockebene erklärt auch ein häufiges Leistungsproblem: False Sharing.

Den Heimatknoten eines Blocks fragen

Jeder Block besitzt einen Eintrag im Heimatverzeichnis. Dieser hält fest, ob der Block nicht im Cache, gemeinsam genutzt oder verändert ist und welche Caches ihn besitzen. Heimatknoten sind auf Knoten oder Verzeichnisteile verteilt; so verteilen sich Speicherbedarf und Anfragebearbeitung.

Ein Lesefehlzugriff geht an den Heimatknoten. Hat der Speicher die aktuellen Daten, kann er sie liefern, und das Verzeichnis erfasst den neuen Mitnutzer. Besitzt ein anderer Cache eine Modified-Kopie, leitet der Heimatknoten die Anfrage an diesen Besitzer weiter, damit aktuelle Daten geliefert werden können. In einem grundlegenden Verzeichnisprotokoll wird dadurch auch der Speicher wieder aktuell.

Um ohne bestehenden Besitz im Zustand Modified zu schreiben, fragt ein Kern den Heimatknoten nach Berechtigung. Das Verzeichnis sendet Invalidierungen an erfasste Mitnutzer, wartet auf Bestätigungen und gewährt exklusiven Besitz. Wiederholte Schreibtreffer in M benötigen dann keine weiteren Kohärenznachrichten. Der Umweg über das Verzeichnis erhöht die Latenz, aber Punkt-zu-Punkt-Verkehr vermeidet es, jeden Cache nach jedem Block zu fragen. Snoop-Filter auf dem Chip erfüllen einen verwandten Zweck, indem sie Anfragen auf mögliche Besitzer beschränken.

Speicherbedarf des Verzeichnisses abschätzen

Ein vollständiger Präsenzvektor enthält ein Bit pro Kern und Block. Bei 64 Kernen sind das 64 Bit oder acht Byte. Im Verhältnis zu einem 64 Byte großen Datenblock kosten allein die Präsenzbits $8\mathrm{/}64=\mathrm{0,125}$, also 12,5 %. Bei 256 Kernen kostet der 32 Byte große Vektor 50 %. Zustandsbits kommen hinzu. Begrenzte Mitnutzerlisten, gröbere Vektoren und die Erfassung nur zwischengespeicherter Blöcke verringern diesen Speicherbedarf.

False Sharing erkennen

True Sharing tritt auf, wenn Kerne Daten über dasselbe Wort austauschen. False Sharing tritt auf, wenn sie auf verschiedene Variablen innerhalb eines Cacheblocks zugreifen und mindestens einer schreibt. Die Variablen sind im Programm unabhängig, aber die Kohärenzeinheit ist größer als jede der beiden Variablen.

Zwei Threads erhöhen getrennte 8 Byte große Zähler an 0x1000 und 0x1008. Bei 64 Byte großen Blöcken sind ihre Blocknummern $\lfloor 4096\mathrm{/}64\rfloor =64$ und $\lfloor 4104\mathrm{/}64\rfloor =64$. Jeder Besitzwechsel überträgt die Berechtigung für beide Zähler zusammen. Häufige abwechselnde Schreibzugriffe können den Block zwischen Caches hin- und herwandern lassen, obwohl keiner der Threads den Wert des anderen benötigt.

Eine Verschiebung des zweiten Zählers nach 0x1040 legt ihn in Block 65. Padding oder Ausrichtung kann Daten pro Thread trennen; auch die Gruppierung von Werten nach ihrem schreibenden Thread hilft. Ein lokaler Akkumulator, der einmal zurückgeschrieben wird, verringert den Verkehr für Besitzwechsel, ohne jeden Wert auffüllen zu müssen.

Auch ein schreibender Kern und ein Leser einer benachbarten Variablen können von False Sharing betroffen sein: Der schreibende Kern invalidiert den ansonsten nutzbaren Block des Lesers. Zwei Leser allein verursachen es nicht. False Sharing erhöht Verkehr und Verzögerung; es führt nicht selbst dazu, dass Aktualisierungen getrennter Variablen falsche Werte erzeugen. Miss Zugriffsmuster und die Platzierung von Cachezeilen, bevor du eine Datenanordnung änderst.

Ein Cluster verbindet vollständige Rechner mit jeweils eigenem Speicher und Betriebssystem. Die Knoten arbeiten durch explizite Nachrichten zusammen. Datenplatzierung und Kommunikationskosten werden daher Teil des Leistungsmodells des Programms.

Daten aufteilen und Ergebnisse zusammenführen

MPI, das Message Passing Interface, gibt jedem Prozess einen Rang und stellt Operationen wie Senden, Empfangen, Rundsenden und Reduzieren bereit. Ein SPMD-Programm kann anhand der Ränge die Daten jedes Prozesses auswählen. Ein Zeiger im privaten Adressraum eines Knotens ist kein gewöhnlicher Verweis über gemeinsamen Speicher auf die Daten eines anderen Knotens.

Sind eine Million Werte bereits auf 100 Knoten verteilt, summiert jeder Knoten lokal 10.000. Eine Baumreduktion führt die Teilsummen in sieben Runden zusammen, denn $\lceil {\log }_{2}100\rceil =7$. Diese Runden umfassen Netzwerknachrichten und nicht nur lokale Additionen. Werden Daten nahe bei der Arbeit gehalten, die sie verarbeitet, sinkt der Übertragungsbedarf.

Cluster skalieren mit Standardkomponenten, und ausgefallene Hardware kann ersetzt werden. Eine Anwendung benötigt nach dem Ausfall eines Knotens möglicherweise dennoch eine Wiederherstellung oder einen Neustart; austauschbare Knoten allein erhalten den Fortschritt einer unterbrochenen Aufgabe nicht.

Mehr als die Verbindungsgeschwindigkeit messen

Latenz ist die Zeit, bis eine Nachricht eintrifft. Verbindungsbandbreite ist die Datenrate einer einzelnen Verbindung. Die Gesamtbandbreite des Netzwerks summiert die Verbindungskapazitäten und beschreibt den idealen Gesamtwert bei Auslastung aller Verbindungen. Die Bisektionsbandbreite summiert die Kapazität über den Schnitt mit der geringsten Kapazität, der die Knoten in gleich große Hälften teilt.

Vergleiche für Verbindungen gleicher Kapazität mit der Bandbreite $B$ Netzwerke mit 64 Knoten anhand einer einheitlichen Zählweise für Verbindungen:

Das Gitter hat acht Reihen mit je sieben horizontalen Verbindungen und acht Spalten mit je sieben vertikalen Verbindungen, also 112. Verbindungen zwischen den gegenüberliegenden Rändern beider Dimensionen ergeben die 128 des Torus. Jeder Hyperwürfelknoten hat sechs Nachbarn, was $64(6)\mathrm{/}2=192$ Verbindungen ergibt, weil jede Verbindung an zwei Endpunkten gezählt wird.

Den Schnitt verfolgen und dann die Kosten betrachten

Ein in zwei Hälften geteilter Ring wird von nur zwei Verbindungen geschnitten. Viele Nachrichten können daher an dieser Stelle konkurrieren, obwohl der Ring insgesamt 64 Verbindungen hat. Ein vollständig verbundenes Netzwerk hat $32(32)=1024$ Verbindungen zwischen seinen beiden Hälften, benötigt aber insgesamt $64(63)\mathrm{/}2=2016$ Verbindungen.

Mehr Verbindungen können die Bandbreite verbessern und die Zahl der Zwischenstationen senken, erhöhen aber den Bedarf an Switch-Ports, Kabeln und Leistungsaufnahme. Gemeinsame Busse haben eine einzige gemeinsam genutzte Kommunikationsressource; Gitter und Tori verteilen Verbindungen über den Raum. Große Cluster verwenden oft mehrstufige vermittelte Netzwerke wie Fat Trees. Ihr Nutzen hängt ebenso vom Verkehrsmuster wie von ihrer Gesamtkapazität ab. Für die Kommunikation zwischen Rechnerhälften zeigt die Bisektion eine Grenze, die die Gesamtbandbreite verbirgt.

Ein Warehouse-Scale Computer verbindet Zehntausende Server mit Stromversorgung, Kühlung und einem Netzwerk, um Internetdienste bereitzustellen. Unabhängige Anfragen liefern parallele Arbeit, während Stapelverarbeitungs-Frameworks wie MapReduce große Datenverarbeitungsaufgaben auf Rechner verteilen.

Den Entwurf auf den Dienst abstimmen

Warehouse-Scale-Systeme legen Wert auf Kosten pro Anfrage, Energieeffizienz und Verfügbarkeit. Die Software wiederholt Arbeit, repliziert Daten und leitet Anfragen an ausgefallenen Rechnern vorbei. Ein Supercomputer konzentriert sich häufiger auf ein einzelnes eng gekoppeltes wissenschaftliches Programm und ein schnelles Verbindungsnetz.

Server sind mit Switches an der Spitze ihrer Racks verbunden, die wiederum über höhere Switch-Ebenen verbunden sind. Zusätzliche Ebenen erhöhen die Kommunikationslatenz. Aufwärtsführende Verbindungen sind oft überbucht, ihre Kapazität liegt also unter der Gesamtkapazität der Serververbindungen. Manche nichtblockierenden Netzwerke nach dem Clos-Prinzip bieten auf den oberen Ebenen genügend Kapazität, um diese Einschränkung zu vermeiden.

Vierzig Server mit Verbindungen von je 10 Gbit/s bieten einem Rack-Switch 400 Gbit/s. Vier aufwärtsführende Verbindungen mit je 25 Gbit/s stellen 100 Gbit/s bereit, das Verhältnis beträgt also 4:1. Senden alle Server gleichzeitig aus dem Rack und teilen sich die Kapazität gleichmäßig, erreicht jeder durchschnittlich 2,5 Gbit/s. Ein gering belasteter Server kann mehr nutzen, wenn andere untätig sind.

Anlagenaufwand von nutzbarer Berechnung trennen

Die Effektivität des Energieeinsatzes, PUE, ist die Gesamtenergie der Anlage geteilt durch die Energie der IT-Geräte im selben Zeitraum. Zehn MW gleichbleibende IT-Last und 1,5 MW für Kühlung und Verteilungsverluste ergeben $PUE=\mathrm{11,5}\mathrm{/}10=\mathrm{1,15}$. Der Idealwert ist 1,0. Google meldete für 2025 einen jährlichen PUE-Wert von 1,09 über alle seine Rechenzentren hinweg.

PUE misst den Anlagenaufwand und nicht die nutzbaren Berechnungen pro Joule. Untätige oder ineffiziente Server können innerhalb der IT Energie verschwenden, selbst wenn die Anlage einen ausgezeichneten PUE-Wert hat.

Bei 50.000 Servern mit durchschnittlich 300 W beträgt die IT-Leistungsaufnahme 15 MW. Bei PUE 1,2 beträgt die Leistungsaufnahme der Anlage 18 MW. Über 8.760 Stunden verbraucht sie damit 157.680 MWh; allein der Zusatzaufwand verbraucht 26.280 MWh. Stromversorgungsanlagen, Kühlung, Strom, Server und Gebäude tragen alle zu den Betriebskosten bei.

Geringe Last kann dennoch viel Leistung aufnehmen

Energieproportionalität bedeutet, dass die Leistungsaufnahme proportional zur Last steigen sollte. Ein Server mit 150 W im Leerlauf und 300 W Spitzenleistung, dessen Leistungsaufnahme linear steigt, benötigt bei 30 % Last $150+\mathrm{0,3}(150)=195$ W. Er benötigt 65 % der Spitzenleistung für 30 % der maximalen Arbeit. Frühere Messungen an Warehouse-Scale-Systemen ergaben, dass viele Server bei etwa 10 % bis 50 % Auslastung arbeiteten. Deshalb ist die Leistungsaufnahme im Leerlauf wichtig.

Ausfälle und langsame Antworten einplanen

Fallen 50.000 Server jeweils durchschnittlich einmal in zehn Jahren aus, sind im Rechnerverbund 5.000 Ausfälle pro Jahr zu erwarten, etwa 14 pro Tag. Diese Erwartung begründet Redundanz und erneute Versuche; sie legt nicht genau 14 Ausfälle für jeden Tag fest.

Eine Anfrage, die Antworten vieler Server benötigt, wartet auf die langsamste erforderliche Antwort. Ist jeder von 100 Servern unabhängig von den anderen mit einer Wahrscheinlichkeit von 1 % langsam, beträgt die Wahrscheinlichkeit, dass keiner langsam ist, ${\mathrm{0,99}}^{100}$. Damit ist mindestens einer mit der Wahrscheinlichkeit $1-{\mathrm{0,99}}^{100}\approx \mathrm{0,634}$ langsam. Die Tail-Latenz eines Servers wird bei großer Auffächerung zu einer häufigen Verzögerung der Anfrage.