Parallelle processors en multicoreprocessors

Parallelle computers voeren bewerkingen tegelijk uit. Ze kunnen de latentie, de tijd die één taak kost, verlagen of de doorvoer, het aantal taken dat per seconde wordt afgerond, verhogen. Die doelen leiden tot verschillende manieren om werk te verdelen.

De vermogensgrens

Door de klokfrequentie te verhogen werd vroeger elke generatie processors sneller. Het dynamisch vermogen volgt ongeveer $P\approx C{V}^{2}f$: capaciteit maal spanning in het kwadraat maal schakelfrequentie. De spanning telt twee keer mee in dat product. Als je zowel de spanning als de frequentie met 15% verlaagt, krijg je ${\mathrm{0,85}}^2\times \mathrm{0,85}=\mathrm{0,614125}$ van het oorspronkelijke dynamische vermogen, ongeveer 61%.

De spanning kon niet voldoende blijven dalen om hogere klokfrequenties te compenseren. Lekstroom en warmte maakten een verdere groei van de frequentie te kostbaar. Deze vermogensgrens dreef de industrie halverwege de jaren 2000 naar multicorechips. Een multicorechip plaatst meerdere processorkernen op één chip; een multiprocessorsysteem heeft minstens twee processors. Eenvoudigere kernen kunnen samen meer werk per joule leveren, mits er genoeg werk te verdelen is.

Vier plaatsen om parallel werk te vinden

Hardware en compilers vinden vooral ILP. Pipelining, meervoudige instructie-uitgifte en uitvoering buiten de oorspronkelijke volgorde laten instructies overlappen terwijl de vereiste resultaten van het programma behouden blijven. Afhankelijkheden en sprongen beperken de mogelijke overlap.

DLP ontstaat als veel elementen dezelfde berekening nodig hebben. SIMD-instructies, vectorprocessors en GPU's benutten deze regelmaat. TLP vereist dat code, een bibliotheek of een compiler het werk over threads verdeelt. Een thread heeft een eigen programmateller, registers en stack; threads binnen één proces delen de adresruimte ervan. Afzonderlijke processen hebben afzonderlijke adresruimtes.

Werk op verzoekniveau vereist weinig samenwerking tussen taken. Extra servers kunnen de capaciteit daardoor bijna evenredig met de toegevoegde middelen verhogen. Ook meerdere onafhankelijke programma's op een laptop gebruiken anders ongebruikte processorkernen en verhogen de totale doorvoer, zelfs als geen van die programma's intern is geparallelliseerd.

Concurrente taken en gelijktijdige uitvoering

Concurrentie betekent dat taken gedurende overlappende perioden voortgang maken. Eén processorkern kan dit doen door de taken om beurten uit te voeren. Parallellisme betekent dat bewerkingen daadwerkelijk op hetzelfde moment op afzonderlijke hardware worden uitgevoerd. Een systeem met één kern kan meerdere programma's tegelijk bezig houden zonder hun instructies gelijktijdig uit te voeren.

Een programma met één sequentiële thread verdeelt zichzelf niet over acht processorkernen alleen omdat die er zijn. De extra kernen kunnen ander werk uitvoeren, maar om dit programma sneller af te ronden, is parallel werk binnen het programma nodig. Ook nadat het werk is verdeeld, beperkt een resterend sequentieel deel de winst.

Een processor kan veel rekeneenheden bevatten zonder veel onafhankelijke instructiestromen te volgen. De taxonomie van Flynn deelt hardware in naar instructiestromen en gegevensstromen. De eerste stap is dus kiezen welk niveau van de machine je beschrijft.

De vier hardwareklassen

Een conventionele processor met één kern is SISD, ook als zijn pipeline meerdere instructies laat overlappen. Superscalaire uitgifte levert op zichzelf geen onafhankelijke programmatellers op en maakt van de kern geen MIMD-hardware.

Een SIMD-instructie kan acht paren 32-bits-drijvendekommagetallen in een 256-bitsregister optellen. De instructie heeft acht rekenbanen, maar die volgen één instructie. SIMD-uitbreidingen met vaste breedte zijn onder meer x86 SSE en AVX en Arm Neon; ook vectorarchitecturen vallen onder de SIMD-klasse. Regelmatige arrayberekeningen passen goed, omdat veel elementen dezelfde bewerking nodig hebben.

MIMD-processors kiezen hun instructies onafhankelijk. Multicorechips, multiprocessors met gedeeld geheugen en clusters vallen in deze klasse. De processorkernen kunnen verschillende programma's uitvoeren, of kopieën van hetzelfde programma die op verschillende momenten verschillende sprongen bereiken.

MISD maakt de 2 × 2-indeling compleet. Er zijn weinig duidelijke voorbeelden in de praktijk. Sommige fouttolerante systemen en verwerkingspipelines worden zo beschreven, maar hun indeling hangt af van hoe hun stromen zijn gedefinieerd.

Eén programma kan toch veel instructiestromen hebben

SPMD, één programma met meerdere gegevensstromen, is een programmeerstijl. Elke werker voert een kopie van hetzelfde programma uit en gebruikt zijn identificatie om zijn aandeel van de gegevens te kiezen. Werkers hebben een onafhankelijke uitvoeringstoestand, dus SPMD draait vaak op MIMD-hardware.

Bij vier processorkernen die 4.000 waarden optellen, verwerkt kern $i$ de indexen $1000i$ tot en met $1000i+999$. Kern 0 verwerkt 0–999; kern 3 verwerkt 3000–3999. Elke kern levert een deelsom en die vier sommen worden gecombineerd. De gedeelde programmatekst vereist niet dat alle vier de kernen in dezelfde cyclus dezelfde instructie uitgeven.

Deel elk niveau afzonderlijk in

Neem een CPU met acht processorkernen waarvan de threads AVX-instructies gebruiken. Over de kernen heen is die MIMD: de kernen gaan onafhankelijk verder. Binnen een AVX-bewerking is die SIMD: één instructie bestuurt meerdere banen. Als je alle rekeneenheden samen zou tellen, zou je deze structuur niet zien.

Ook een cluster dat kopieën van één wetenschappelijk programma uitvoert, is MIMD-hardware met SPMD-software. De knooppunten communiceren via berichten. Een GPU combineert veel onafhankelijk geplande threadgroepen met SIMD-achtige uitvoering binnen elke groep; NVIDIA noemt zijn uitvoeringsstijl op basis van threads SIMT.

Om een beschrijving in te delen, vraag je wie de volgende instructie kiest. Vraag daarna of die instructie één gegevenselement of meerdere verwerkt. Het aantal bronbestanden, geïnstalleerde programma's of rekeneenheden alleen kan die twee vragen niet beantwoorden.

Een cachemisser kan een processorkern honderden cycli laten wachten. Hardwarematige multithreading bewaart de uitvoeringstoestand van meerdere threads binnen één kern, zodat tijdens die wachttijd nuttige instructies van een andere thread kunnen worden uitgevoerd.

Wat wordt verdubbeld en wat wordt gedeeld

Elke hardwarethread heeft een eigen programmateller en registertoestand. De threads delen de uitvoeringseenheden en caches van de processorkern en de toegang tot geheugenbandbreedte. Door de toestand in hardware te bewaren gaat wisselen veel sneller dan het opslaan en herstellen van een proces door het besturingssysteem, wat duizenden cycli kan kosten. Hardware kan in één cyclus of minder wisselen.

Het besturingssysteem ziet elke hardwarethread als een logische processor. Acht fysieke processorkernen met twee hardwarethreads per kern tonen 16 logische processors. Dat aantal beschrijft de beschikbare threadcontexten, niet 16 volledige fysieke kernen.

Drie planningsaanpakken

Fijnmazige planning kan tussen gereedstaande threads wisselen, ook als de vorige thread verder zou kunnen. Grofmazige planning laat die thread draaien totdat een kostbare wachttijd een wisseling rechtvaardigt. SMT gebruikt een processorkern met meervoudige instructie-uitgifte en dynamische planning: instructies van verschillende threads kunnen tijdens één cyclus verschillende vrije uitgifteplaatsen bezetten. Registerhernoeming houdt hun toestand gescheiden. Hyper-Threading van Intel is een SMT-implementatie; sommige IBM POWER-processors ondersteunen acht hardwarethreads per kern.

Reken het verbergen van latentie door

Stel dat een thread 20 cycli rekent en daarna 200 cycli op het geheugen wacht. De cyclus van de thread is 220 cycli lang. Alleen neemt nuttig werk $20\mathrm{/}220\approx \mathrm{0,091}$, oftewel 9,1% van de tijd in beslag.

Met vier threads, vrij gespreid werk en zonder wisselkosten of extra concurrentie om middelen kan nuttig werk $4\times 20\mathrm{/}220\approx \mathrm{0,364}$, oftewel 36,4%, in beslag nemen. De processorkern heeft nog steeds tussenpozen, omdat vier rekenperioden van 20 cycli de wachttijden niet kunnen overbruggen.

Elf threads leveren samen 220 cycli aan rekenperioden. Terwijl één thread 200 cycli wacht, kunnen de andere tien elk 20 cycli rekenen. Deze ideale planning houdt de kern bezig. De werkelijke prestaties hangen ook af van de vraag of die threads concurreren om geheugen en uitvoeringsmiddelen.

Doorvoer en de voltooiingstijd van een thread

Multithreading verhoogt vooral de doorvoer. Een bepaalde thread kan langer duren doordat andere threads cycli en cacheruimte gebruiken die de thread anders zou krijgen. Twee hardwarethreads die één rekeneenheid delen, kunnen niet allebei de capaciteit van een tweede volledige processorkern erbij krijgen.

Als beide threads al dezelfde volledig bezette uitvoeringseenheden nodig hebben, heeft SMT weinig vrije capaciteit om te vullen. Als één thread wacht terwijl de andere kan worden uitgevoerd, is delen nuttig. GPU's passen het verbergen van latentie op grotere schaal toe door veel threadgroepen gereed te houden voor hun planners.

Een parallelle machine is nuttig als die de verstreken tijd verandert. De versnellingsfactor vergelijkt de beste sequentiële uitvoering met de parallelle uitvoering; de efficiëntie vergelijkt die winst met het aantal gebruikte processors.

Begin met de gemeten tijd

Voor $p$ processors geldt $S(p)=T_1\mathrm{/}{T}_p$ en $E(p)=S(p)\mathrm{/}p$. Een sequentiële taak van 100 s die op acht processorkernen in 16 s klaar is, heeft een versnellingsfactor van $100\mathrm{/}16=\mathrm{6,25}$. De efficiëntie is $\mathrm{6,25}\mathrm{/}8=\mathrm{0,78125}$, ongeveer 78%. Meet voor beide uitvoeringen de verstreken tijd en gebruik dezelfde vaste taak.

Verdeel de tijd in een verbeterd en een onveranderd deel

In de wet van Amdahl is $f$ de fractie van de oorspronkelijke uitvoeringstijd die een verbetering versnelt. Het is een fractie van de tijd, geen fractie van de regels broncode. Als dat deel $s$ keer zo snel wordt, geldt

Het onveranderde deel kost nog steeds de volledige oorspronkelijke tijd. Voor een ideale parallelle verdeling over $p$ processors stel je $s=p$. Dit model laat overhead door communicatie, ongelijke werklastverdeling en synchronisatie buiten beschouwing; die kosten kunnen de gemeten tijd verlengen.

Bij 90% parallelle tijd en tien processors is de genormaliseerde tijd $\mathrm{0,1}+\mathrm{0,9}\mathrm{/}10=\mathrm{0,19}$. De versnellingsfactor is $1\mathrm{/}\mathrm{0,19}\approx \mathrm{5,26}$ en de efficiëntie is ongeveer 53%. Met 100 processors is de tijd $\mathrm{0,1}+\mathrm{0,009}=\mathrm{0,109}$ en de versnellingsfactor ongeveer 9,17. Een vertienvoudiging van het aantal processors heeft de versnellingsfactor niet verdubbeld.

Bepaal de grens voordat je meer processors koopt

Als het aantal processors groeit, nadert de parallelle tijd tot nul. De sequentiële fractie blijft bestaan en geeft de bovengrens $S_{max}=1\mathrm{/}(1-f)$. Een sequentiële fractie van 10% begrenst de ideale versnellingsfactor op 10; een fractie van 5% op 20. Eindige aantallen processors naderen deze bovengrenzen van onderaf.

De wet van Amdahl beoordeelt ook één snellere instructieklasse. Als vermenigvuldigen 30% van een uitvoering kost en vier keer zo snel wordt, wordt de genormaliseerde tijd $\mathrm{0,7}+\mathrm{0,3}\mathrm{/}4=\mathrm{0,775}$. Het hele programma is slechts $1\mathrm{/}\mathrm{0,775}\approx \mathrm{1,29}$ keer zo snel.

Reken terug vanaf een doel

Om een versnellingsfactor van 90 op 100 processors te behalen, schrijf je $1\mathrm{/}90=(1-f)+f\mathrm{/}100=1-\mathrm{0,99}f$. Dus $f=(1-1\mathrm{/}90)\mathrm{/}\mathrm{0,99}\approx \mathrm{0,99888}$. Slechts ongeveer 0,11% van de oorspronkelijke tijd kan sequentieel blijven. Een doel dicht bij het aantal processors vereist een zeer kleine onveranderde fractie.

Een gemeten versnellingsfactor kan soms groter zijn dan het aantal processors. Door een grote werkverzameling op te splitsen kan elk deel in de cache passen, waardoor de kosten per bewerking dalen. Deze superlineaire versnellingsfactor verandert een uitgangspunt van het eenvoudige model: parallel werk is niet langer slechts het oorspronkelijke werk verdeeld in gelijke delen. Controleer het algoritme, de gegevensomvang en het cachegedrag voordat je zo'n meting interpreteert.

Extra processors kunnen dezelfde taak verkorten of een machine een grotere taak laten verwerken. Sterke schaling en zwakke schaling meten deze verschillende toepassingen en beide hangen af van hoeveel nuttig werk overblijft na de coördinatiekosten.

Vast werk en groeiend werk

Bij sterke schaling blijft de probleemomvang vast. Bij zwakke schaling groeit het werk evenredig met het aantal processors, zodat het werk per processor gelijk blijft. Bij perfecte zwakke schaling blijft de verstreken tijd gelijk terwijl het totale werk groeit.

Neem tien sequentiële optellingen gevolgd door de elementsgewijze optelling van twee matrices. Een 10 × 10-matrix vereist 100 parallelle optellingen. De totale tijd is 110 eenheden op één processor, 20 op tien en 12 op vijftig: versnellingsfactoren van 5,5 en ongeveer 9,2. Een 40 × 40-matrix vereist 1.600 parallelle optellingen. De tijden worden 1.610, 170 en 42 eenheden: versnellingsfactoren van ongeveer 9,5 en 38,3. Bij het grotere probleem weegt nuttig parallel werk zwaarder ten opzichte van dezelfde tien sequentiële optellingen.

Het andere uitgangspunt van Gustafson

De wet van Gustafson gebruikt de sequentiële tijdsfractie die tijdens de parallelle uitvoering is gemeten. Als je die fractie $s$ noemt, is de geschaalde versnellingsfactor $S_G=s+(1-s)p$. Voor $s=\mathrm{0,05}$ en 100 processors geldt $S_G=\mathrm{0,05}+\mathrm{0,95}(100)=\mathrm{95,05}$.

De wet van Amdahl met 5% sequentiële tijd tijdens de oorspronkelijke uitvoering op één processor geeft op 100 processors slechts ongeveer 16,8. Deze fracties gebruiken verschillende noemers en beantwoorden verschillende vragen. Een geschaalde versnellingsfactor van 95,05 betekent niet dat de oorspronkelijke taak met vaste omvang 95,05 keer zo snel klaar was.

De langzaamste werker bepaalt de eindtijd

Als 400 gelijke werkeenheden naar acht processorkernen gaan, is het evenwichtige aandeel 50 per kern. Stel dat één kern er 120 krijgt en de andere zeven elk 40. De stap duurt 120 tijdseenheden, dus de versnellingsfactor is $400\mathrm{/}120\approx \mathrm{3,33}$. Als je door de gemiddelde werklast zou delen, zou je het wachten op de drukste kern missen.

Ook het aanmaken en plannen van threads, berichten, barrières en concurrentie om gedeelde bandbreedte kosten tijd. Korrelgrootte is de hoeveelheid rekenwerk tussen communicatie- of synchronisatiemomenten. Grotere stukken verminderen de coördinatiekosten per nuttige bewerking, maar bij te weinig stukken kunnen sommige werkers niets te doen hebben.

Meer processors kunnen de tijd verlengen

Voor een volledig parallelle taak van 100 s met $\mathrm{0,5}p$ seconden overhead geldt $T_p=100\mathrm{/}p+\mathrm{0,5}p$. Tien processors doen er 15 s over. Veertien doen er ongeveer 14,14 s over, dicht bij het minimum rond $\sqrt{200}$ processors. Veertig doen er 22,5 s over, omdat hun 20 s overhead zwaarder weegt dan de kortere nuttige berekening.

Vermeld bij het rapporteren van een resultaat of het werk gelijk bleef of groeide en vergelijk met de tijd van het beste sequentiële programma. Als je een trage parallelle implementatie op één processor als uitgangspunt gebruikt, kan de gerapporteerde versnellingsfactor te hoog uitvallen.

De opgegeven rekensnelheid van een processor is slechts één grens. Het rooflinemodel vraagt ook of het geheugen snel genoeg gegevens kan leveren om die snelheid vol te houden. Het model geeft een bovengrens, die helpt om een nuttige richting voor optimalisatie te bepalen.

Tel bewerkingen per DRAM-byte

Rekenintensiteit is het aantal drijvendekommabewerkingen gedeeld door het aantal bytes dat tussen DRAM en de processor wordt overgedragen. Het oorspronkelijke werk over het rooflinemodel noemt deze op DRAM gebaseerde maat operationele intensiteit. Tel het verkeer op die grens, inclusief lees- en schrijfbewerkingen.

DAXPY berekent $y_i=a{x}_i+y_i$ met getallen met dubbele precisie. Elk element vereist één vermenigvuldiging en één optelling: twee FLOP. Een waarde van 8 bytes uit elke array lezen en 8 bytes terugschrijven verplaatst 24 bytes. De intensiteit is $I=2\mathrm{/}24\approx \mathrm{0,0833}$ FLOP per byte. De scalaire coëfficiënt wordt hergebruikt; in dit verkeersmodel wordt die niet voor elk element opnieuw uit DRAM geladen.

Pas beide grenzen toe

Als de piekrekencapaciteit $P$ is en de piekgeheugenbandbreedte $B$, is het plafond $\min (P,BI)$. Bandbreedte maal FLOP per byte heeft als eenheid FLOP per seconde. In een dubbel-logaritmische grafiek loopt de bandbreedtegrens omhoog en is de rekengrens horizontaal.

Bij 1.000 GFLOP/s en 100 GB/s is het knikpunt van het rooflinemodel $P\mathrm{/}B=1000\mathrm{/}100=10$ FLOP per byte. Onder deze intensiteit ligt het bandbreedteplafond lager; erboven ligt het rekenplafond lager. DAXPY wordt begrensd door $100(2\mathrm{/}24)\approx \mathrm{8,33}$ GFLOP/s, minder dan 1% van de piekrekencapaciteit van de machine. Een kernel met intensiteit 2 wordt begrensd door 200 GFLOP/s.

Een door geheugenbandbreedte begrensde kernel kan nog steeds profiteren van extra processorkernen of bredere SIMD zolang er te weinig verzoeken tegelijk onderweg zijn om de beschikbare bandbreedte te benutten. Zodra die bandbreedte volledig wordt benut, verhoogt extra rekencapaciteit het plafond niet. Meer bandbreedte of meer nuttig rekenwerk per overgedragen byte doet dat wel.

Hergebruik verandert de intensiteit

Het vermenigvuldigen van twee 1000 × 1000-matrices met dubbele precisie voert ongeveer $2(1000{)}^3=2\times {10}^9$ FLOP uit. Als elk element van de twee invoermatrices en de uitvoermatrix één keer de DRAM-grens passeert, is het totale verkeer 24 MB, wat ongeveer 83 FLOP per byte geeft. Dat ligt boven het knikpunt van de voorbeeldmachine.

Een naïeve implementatie kan matrixwaarden herhaaldelijk uit DRAM lezen. Blokgewijze verwerking organiseert lussen rond deelmatrices die in de cache passen, zodat waarden voor veel bewerkingen worden gebruikt voordat ze worden vervangen. Het algoritme blijft matrixvermenigvuldiging, maar de implementatie verplaatst minder DRAM-bytes per bewerking.

Een plafond is geen gemeten resultaat

Een kernel kan onder de bovengrens presteren door afhankelijkheden, instructieoverhead, ongunstige toegangspatronen of onvoldoende parallel werk. Lagere prestatieplafonds tonen zulke beperkingen: gebrek aan SIMD kan het rekenplafond verlagen, terwijl ongunstige toegangspatronen de bruikbare bandbreedte verlagen. Een punt onder een plafond vraagt om onderzoek; het plafond alleen voorspelt niet de exacte verwerkingssnelheid.

Een SIMD-uitbreiding verdeelt een breed register in banen. Eén instructie voert in elke baan dezelfde bewerking uit, waardoor regelmatige lussen over onafhankelijke elementen goed passen.

Tel banen en iteraties

Het aantal banen is de registerbreedte gedeeld door de elementbreedte. Een 256-bitsregister bevat acht 32-bits-drijvendekommagetallen met enkele precisie, vier 64-bitsgetallen met dubbele precisie of 32 elementen van één byte. SSE gebruikt 128-bits-XMM-registers; AVX en AVX2 gebruiken 256-bits-YMM-registers. In 64-bitsmodus zijn er 16 YMM-registers. AVX-512 biedt 32 512-bits-ZMM-registers, die elk 16 drijvendekommagetallen met enkele precisie bevatten. Arm Neon gebruikt 128-bitsregisters.

Voor c[i] = a[i] + b[i] over 1.000 drijvendekommagetallen verwerken instructies met acht banen de array in 125 vectoriteraties. Elke iteratie laadt twee vectoren, telt hun banen op en slaat een vector op. Bij 1.003 elementen blijven na 125 volledige iteraties drie elementen over. Een scalaire restlus of een bewerking met een masker handelt dat korte restant af zonder buiten de arrays te komen.

Acht banen vormen een rekencapaciteit, geen garantie dat het programma acht keer zo snel wordt. Voorbereiding, het afhandelen van de rest, sequentieel werk en geheugenbandbreedte kosten nog steeds tijd.

Kies een indeling die een register vult

Eenvoudige vectorlaadinstructies werken goed op aaneengesloten elementen met stapgrootte één. Een array van {x, y, z}-punten met waarden van 4 bytes plaatst opeenvolgende x-velden 12 bytes uit elkaar. Een structuur van arrays plaatst de x-waarden bij elkaar, daarna de y-waarden en vervolgens de z-waarden. Als je alleen x verwerkt, kun je nu met één aaneengesloten laadinstructie een vector vullen. Verzamel- en verspreidinstructies ondersteunen verspreide adressen, maar meestal tegen hogere toegangskosten.

Houd voorwaardelijk werk binnen banen

Voor c[i] = (a[i] > 0) ? a[i] : 0 maakt een vergelijking een masker per baan. Een selectie-instructie kiest tussen kandidaatresultaten, of instructies met een masker werken alleen de geselecteerde banen bij. Inactieve banen behouden of ontvangen waarden volgens de maskerwerking van de instructie.

Als verschillende banen lange, verschillende takken nodig hebben, moeten de instructies voor die paden nog steeds worden uitgegeven, met alleen de relevante banen actief. De groep kan niet elke baan een afzonderlijke instructiestroom geven. Een masker bewaart de correctheid, maar kan niet al die padinstructies laten verdwijnen.

Controleer de onafhankelijkheid van iteraties

Compilers kunnen lussen automatisch vectoriseren; programmeurs kunnen ook bibliotheken of intrinsieke functies zoals _mm256_add_ps gebruiken. De compiler moet vaststellen dat gelijktijdige iteraties het resultaat behouden. De recurrente berekening a[i] = a[i-1] + b[i] hangt af van de nieuwe waarde uit de vorige iteratie en kan niet rechtstreeks als onafhankelijke elementen worden gevectoriseerd. Mogelijke overlap van pointers, aliasing genoemd, kan ook een overlapcontrole tijdens de uitvoering vereisen of vectorisatie verhinderen.

Een reductie bewaart meerdere deelsommen in een register en combineert daarna de banen. Een andere volgorde van drijvendekommaoptellingen kan de laatste bits veranderen, omdat optellen niet associatief is. Sommen van gehele getallen komen overeen als de gekozen rekenwijze geen overloop heeft of gedefinieerde modulaire rekenkunde gebruikt.

SIMD en multicore-uitvoering zijn te combineren. Acht processorkernen die elk één optelinstructie voor acht banen per cyclus uitgeven, kunnen maximaal 64 drijvendekommaoptellingen per cyclus uitvoeren. Ze delen nog steeds grenzen zoals de geheugenbandbreedte.

Een vectorarchitectuur drukt een bewerking op een reeks elementen uit en bewaart de actieve vectorlengte in een register. Hetzelfde programma kan draaien op implementaties met verschillende capaciteiten van vectorregisters.

Onderscheid capaciteit, lengte en banen

Versie 1.0 van de RISC-V-vectoruitbreiding, RVV, is in 2021 bekrachtigd. Die biedt 32 vectorregisters, v0 tot en met v31. Elke implementatie kiest de registerbreedte, VLEN, als een macht van twee. De elementbreedte en de groepering van registers bepalen hoeveel elementen passen.

De actieve lengte, vl, is het aantal elementen dat de volgende vectorinstructies verwerken. vsetvli configureert de vectorbewerking op basis van het aantal resterende elementen en het gewenste elementformaat en geeft de gekozen vl terug. De code gaat met die teruggegeven waarde vooruit. Dit is belangrijk omdat overdraagbare code de werkelijk door de implementatie gekozen lengte moet gebruiken.

Banen zijn de parallelle pipelines die de vectorbewerking uitvoeren. Een optelling van 32 elementen op vier banen duurt na het opstarten van de pipeline ongeveer acht cycli. Een vector van 32 elementen vereist geen 32 rekeneenheden. Langere vectoren en meer banen zijn afzonderlijke ontwerpkeuzes.

Verdeel een lus in vectorblokken

Lusopdeling in vectorblokken verdeelt een lange lus in stukken ter grootte van een vector. Voor 100 elementen en maximaal 32 per bewerking is één geldige reeks lengtes 32, 32, 32 en 4. Elke doorgang laadt het volgende deel, berekent het, slaat het op en vermindert het resterende aantal met de vl van die doorgang.

DAXPY over 1.000 drijvendekommagetallen met dubbele precisie kan op dezelfde manier 31 doorgangen van 32 gebruiken, gevolgd door een doorgang van acht, in totaal 32 doorgangen. RVV kan tegen het einde andere toegestane lengtes kiezen, dus de lus gebruikt de teruggegeven vl in plaats van een bepaalde laatste verdeling aan te nemen. Door de actieve lengte in te stellen wordt het korte restant afgehandeld zonder aparte scalaire restlus.

Het patroon voor de lusbesturing is:

1. Vraag een lengte voor de resterende elementen aan.
2. Laad en bereken dat aantal elementen.
3. Sla de resultaten op.
4. Schuif de pointers op en trek de teruggegeven lengte af.

Eén vectorinstructie beschrijft veel elementbewerkingen. Dat vermindert het ophaal- en decodeerwerk ten opzichte van een scalaire lusinhoud die voor elk element wordt herhaald.

Verwerk afhankelijke bewerkingen en niet-aaneengesloten gegevens

Koppeling van vectorbewerkingen laat een afhankelijke vectorinstructie met de eerste resultaten beginnen voordat de producerende instructie klaar is met latere elementen. Bij een vermenigvuldiging gevolgd door een optelling kan het eerste product de optelpipeline ingaan terwijl latere producten nog worden gevormd.

Een masker, in RVV gewoonlijk in register v0, selecteert welke elementen worden bijgewerkt. Een laadinstructie met vaste stapgrootte gebruikt een vaste afstand in bytes. In een rijgewijs opgeslagen 100 × 100-matrix met drijvendekommagetallen met dubbele precisie ga je bij een stap omlaag in één kolom 100 elementen, oftewel 800 bytes, verder. Geïndexeerde laad- en opslaginstructies gebruiken expliciete adressen voor onregelmatige gegevens.

Zowel vectorarchitecturen als uitbreidingen met vaste breedte zijn SIMD in de taxonomie van Flynn. Uitbreidingen met vaste breedte coderen specifieke breedtes; RVV-code drukt lengtes uit die zich aan de beschikbare hardware aanpassen. Beide hebben nog steeds een uitvoeringsvolgorde nodig die rekening houdt met de gegevensafhankelijkheden van de lus.

Een CPU concentreert middelen op lage latentie voor enkele threads. Een GPU concentreert middelen op hoge doorvoer over veel dataparallelle threads. Om ertussen te kiezen moet je naar de hele taak kijken, inclusief gegevensverplaatsing.

Threads draaien op streamingmultiprocessors

Een GPU van NVIDIA bevat veel streamingmultiprocessors, of SM's. Een SM heeft rekenbanen, warpplanners, een grote registerbank en snel gedeeld geheugen op de chip. Een “CUDA-core” is de naam voor een rekenbaan, geen volledige onafhankelijke kern zoals in een CPU. AMD noemt de overeenkomstige verwerkingsgroep een compute unit.

NVIDIA groepeert threads in warps van 32. Eén uitgegeven instructie bewerkt de actieve threads van de warp, elk met eigen registerwaarden en gegevens. Dit is SIMT, één instructie, meerdere threads. AMD gebruikt wavefronts, met 32 of 64 threads afhankelijk van de architectuur.

Een SM houdt veel warps actief beschikbaar. Als één warp op het geheugen wacht, kiest de planner een andere gereedstaande warp. Meer gereedstaande threads dan banen verbergen latentie; ze verminderen niet de latentie van een afzonderlijke DRAM-toegang.

Divergentie laat banen ongebruikt

Als threads in één warp verschillende takken volgen, voert de warp de vereiste paden uit terwijl de andere threads met een masker zijn uitgeschakeld. Een gelijk verdeelde if-else met even kostbare takken kan ongeveer twee keer de tijd van één pad kosten, omdat beide paden moeten worden uitgevoerd. Een sprong kan onschadelijk zijn als de voorwaarde binnen elke warp gelijk is, zelfs als verschillende warps verschillende paden kiezen.

Laat geheugentoegangen samenwerken

Registers per thread bevatten privéwaarden. Gedeeld geheugen per blok is opslag op de chip die door het programma wordt beheerd. L1- en L2-caches zorgen voor automatisch hergebruik, terwijl globaal geheugen het DRAM van de GPU is, zoals GDDR of HBM.

Als 32 threads opeenvolgende woorden van 4 bytes lezen, kan hun bereik van 128 bytes met een klein aantal brede transacties worden bediend. Deze samenvoeging van geheugentoegangen gebruikt de bandbreedte efficiënt. Verspreide adressen kunnen voor dezelfde nuttige gegevens veel transacties vereisen. Hergebruik via gedeeld geheugen kan herhaald verkeer naar het globale geheugen verder verminderen.

Neem de overdrachtstijd mee

Bij een afzonderlijke GPU gaan invoer en resultaten over een verbinding zoals PCIe of NVLink. 800 MB versturen en 400 MB ontvangen met een effectieve snelheid van 20 GB/s draagt 1,2 GB over in 0,06 s, oftewel 60 ms. Met een kernel van 10 ms erbij is dat in totaal 70 ms. Een CPU die in 50 ms klaar is, wint deze vergelijking ondanks de veel snellere GPU-berekening.

Door gegevens gedurende meerdere kernels op de GPU te houden, voorkom je herhaalde overdrachten. Grote, regelmatige taken kunnen ook genoeg nuttig rekenwerk bieden om de kosten van verplaatsing en starten te overtreffen. Kleine, sequentiële taken met veel sprongen zijn vaak gunstiger voor een CPU.

Specialiseer verder als het werk dat toelaat

Een heterogeen systeem gebruikt de CPU om het programma en het besturingssysteem uit te voeren en een hardwareversneller voor geschikte onderdelen. Domeinspecifieke hardwareversnellers zoals TPU's en NPU's besteden hardware aan een beperktere taak, vaak matrixvermenigvuldiging met 8-bits gehele getallen of 16-bits-drijvendekommagetallen. Specialisatie kan de prestaties per watt in dat domein verbeteren, terwijl het bereik van werk dat de hardware goed verwerkt kleiner wordt.

Een GPU-kernel draait één keer per gestarte thread. De host start een threadblokraster en elke thread berekent welk deel van de gegevens hij verwerkt. CUDA noemt de CPU de host en de GPU het apparaat; OpenCL gebruikt de verwante termen work-item en work-group.

Bereken een unieke index

Bij een eendimensionale start is de elementindex blockIdx.x * blockDim.x + threadIdx.x. Met 256 threads per blok verwerkt thread 5 in blok 2 element 517: de twee voorafgaande blokken zijn goed voor 512 elementen en de thread voegt een offset van vijf toe. Thread- en blokindexen beginnen bij nul.

Een SAXPY-kernel past één vermenigvuldiging en één optelling toe op elk drijvendekommagetal:

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

Het raster vervangt de gewone lus over elementen. Elke actieve thread werkt zijn eigen y[i] bij; aangrenzende threads gebruiken aangrenzende adressen, wat samenvoeging van geheugentoegangen ondersteunt.

Rond het aantal blokken naar boven af

Voor 1.000.000 elementen en 256 threads per blok heeft het raster $\lceil 1000000\mathrm{/}256\rceil =3907$ blokken nodig. Die starten 1.000.192 threads. De laatste 192 hebben indexen buiten de array, dus de voorwaarde i < n voorkomt hun geheugentoegangen. Naar beneden afronden zou te weinig threads starten en geldige elementen ongemoeid laten.

Een blok van 256 threads vormt acht warps. Blokken kunnen aan elke SM worden toegewezen en in elke volgorde worden uitgevoerd. De correctheid van het programma mag er niet van afhangen dat blok 0 vóór blok 1 klaar is.

Organiseer het werk van host en apparaat

Een gebruikelijke host reserveert opslag op het apparaat, kopieert de invoer, start saxpy<<<3907, 256>>>(n, 2.0f, d_x, d_y) en haalt de resultaten op. De start is asynchroon: de CPU kan doorgaan terwijl de kernel draait. De host moet via synchronisatie of een geschikte blokkerende kopieeractie van de resultaten wachten voordat hij de voltooide gegevens gebruikt.

Werk samen binnen een blok

Threads in één blok kunnen waarden uitwisselen via gedeeld geheugen en elkaar bij __syncthreads() treffen. Elke thread die aan die barrière voor het hele blok deelneemt, moet de barrière bereiken. De barrière binnen if (i < n) plaatsen is onveilig als slechts sommige threads aan de voorwaarde voldoen.

Voor een reductie van 1.024 waarden in gedeeld geheugen tellen 512 threads paren op, daarna combineren 256 de resultaten, daarna 128, enzovoort. Na tien halveringsstappen blijft één som over. Een barrière tussen de stappen voorkomt dat een thread een deelresultaat gebruikt voordat de producerende thread het schrijft.

Een gewone start heeft geen barrière voor het hele raster. Afzonderlijke blokken kunnen deelsommen schrijven en een tweede kernel kan die combineren, of geschikte atomaire operaties kunnen een globaal resultaat bijwerken. Een blokbarrière alleen kan verschillende blokken niet coördineren. Een lus met een stapgrootte gebaseerd op het raster is een ander toewijzingspatroon: elke thread verwerkt meerdere elementen en gaat telkens met de totale rastergrootte vooruit.

Een multiprocessor met gedeeld geheugen geeft zijn processors één fysieke adresruimte. Threads communiceren met laad- en opslaginstructies voor gedeelde variabelen, terwijl ze ook privéwaarden zoals lokale deelsommen bewaren.

Gedeelde adressen kunnen verschillende kosten hebben

In UMA, uniforme geheugentoegang, is de geheugentoegangstijd vanaf elke processor ongeveer gelijk. In NUMA, niet-uniforme geheugentoegang, bereikt een processorsocket zijn lokale geheugen sneller dan het geheugen van een andere socket. De adresruimte is gedeeld, hoewel het fysieke geheugen is verdeeld.

Bij 70% lokale toegangen van 80 ns en 30% toegangen op afstand van 140 ns is het gewogen gemiddelde $\mathrm{0,7}(80)+\mathrm{0,3}(140)=98$ ns. Door gegevens dicht bij de belangrijkste gebruiker te plaatsen, verminder je het aandeel toegangen op afstand. Als je elke toegang als lokaal meetelt, onderschat je de tijd.

Een gebruikelijke cache-indeling geeft elke processorkern eigen L1-instructie- en datacaches, vaak een eigen L2-cache en een gedeelde cache op het laatste niveau. Een ring of rooster verbindt de kernen en caches. Gedeeld geheugen beschrijft de adressering; het vereist niet dat elke kern één gemeenschappelijke cache gebruikt.

Verdeel, bereken en reduceer

Om 64.000 waarden op 64 processorkernen op te tellen, verwerkt elke kern 1.000 waarden. De 64 deelsommen kunnen in een reductieboom worden gecombineerd. Het aantal actieve optellers halveert bij elke stap: 32, 16, 8, 4, 2 en 1. Zes stappen combineren de deelresultaten.

Met ongeveer 1.000 lokale optelstappen plus zes combinatiestappen krijg je ongeveer 1.006 stappen langs het langste pad, vóór overhead. De boom voorkomt dat elke invoer één gedeelde teller bijwerkt. Elke stap hangt af van de voltooide waarden uit de vorige stap, dus de werkers komen tussen de stappen samen bij synchronisatiebarrières.

OpenMP kan dit uitdrukken met #pragma omp parallel for reduction(+:sum). Dit verdeelt iteraties, geeft threads eigen deelsommen en combineert die. POSIX-threads bieden explicietere controle over threadbeheer.

Maak een update ondeelbaar

Een atomaire lees-wijzig-schrijfoperatie maakt een update ondeelbaar zoals andere processors die waarnemen. Atomair omwisselen en atomair ophalen en optellen zijn voorbeelden. Een laadinstructie gevolgd door een gewone opslaginstructie zijn twee afzonderlijke operaties, dus een andere thread kan tussendoor ingrijpen.

RISC-V biedt lr.w en sc.w. De laadoperatie met reservering leest een woord en legt een reservering vast. De voorwaardelijke opslagoperatie probeert te schrijven en meldt succes of mislukking. Een schrijfbewerking tussendoor naar de gereserveerde locatie laat de poging mislukken; de software probeert het opnieuw. De reservering blokkeert andere processors niet en de poging kan ook om andere toegestane redenen mislukken. Deze basisoperaties ondersteunen synchronisatiemechanismen.

Let op de gedeelde knelpunten

Als elke processorkern 15 GB/s aan gegevens kan verwerken maar de processorsocket 60 GB/s levert, benutten ongeveer vier van zulke stromen de bandbreedte volledig. Meer kernen kunnen de gegevenssnelheid van de socket niet verhogen. Communicatie via caches kost ook werk: een schrijver zonder schrijfrecht moet eigenaarschap verkrijgen en andere kopieën ongeldig maken. Een schrijver die al een blok in de toestand Modified heeft, kan lokaal doorgaan; latere lezers moeten mogelijk een nieuwe kopie ophalen.

Privécaches verbeteren de latentie, maar kunnen meerdere kopieën van dezelfde geheugenlocatie bevatten. Cachecoherentie voorkomt dat die kopieën tegenstrijdige geschiedenissen van de schrijfbewerkingen naar die locatie opleveren.

Waarom het geheugen bijwerken niet voldoende is

Begin met X = 0 in het geheugen. Processorkern A leest nul en bewaart die waarde in zijn cache; kern B doet hetzelfde. A schrijft één in zijn cache die wijzigingen terugschrijft. Zonder coherentieprotocol kan B een treffer in zijn eigen cache krijgen en nul blijven lezen. Direct doorschrijven naar het geheugen zou de kopie die al in de cache van B staat niet herstellen.

Coherentie werkt op het niveau van cacheblokken, gewoonlijk van 64 bytes. Een controller houdt de toestand van elk blok bij en volgt protocolregels wanneer processors het blok lezen of schrijven. Migratie verplaatst nuttige gegevens naar de cache van de huidige gebruiker. Replicatie laat meerdere lezers schone kopieën bewaren. Coherentie moet beide voordelen behouden en tegelijk schrijfbewerkingen beheersen.

Drie vereiste eigenschappen

Als naar een locatie eerst één en daarna twee wordt geschreven, kunnen lezers het niet oneens zijn over welke schrijfbewerking later kwam. De latere schrijfbewerking waarnemen en daarna de eerdere zou die geschiedenis omkeren. Een nieuwe schrijfbewerking daarna zou natuurlijk opnieuw de waarde één kunnen opslaan; serialisatie gaat over schrijfbewerkingen, niet over de vraag of een getalswaarde ooit kan terugkeren.

Beheer eigenaarschap en kopieën

In een protocol met schrijfinvalidatie verkrijgt een schrijver exclusief schrijfrecht en maakt hij andere kopieën in caches ongeldig. Een latere leesbewerking van een ongeldig gemaakt blok levert een misser op en haalt actuele gegevens op. Dit is een coherentiemisser, naast eerste-toegangs-, capaciteits- en conflictmissers.

Een protocol met schrijfbijwerking stuurt in plaats daarvan nieuwe waarden naar bestaande kopieën. Herhaald schrijven kan dit kostbaar maken. Protocollen die meeluisteren op de bus versturen coherentieverzoeken naar iedereen, zodat caches kunnen reageren. Een coherentieregister houdt bij welke caches kopieën delen en stuurt berichten naar de relevante caches.

Coherentie, volgorde en atomiciteit

Coherentie gaat over één locatie. Een geheugenconsistentiemodel definieert toegestane waarneembare volgordes en resultaten van lees- en schrijfbewerkingen, inclusief operaties op verschillende locaties. Sequentiële consistentie vereist een uitkomst die overeenkomt met een gezamenlijke volgorde die de programmavolgorde van elke processor bewaart. Zwakkere modellen staan meer herschikkingen toe; instructies voor geheugenbarrières dwingen waar nodig de vereiste volgorde af.

Coherente kopieën noch een consistente volgorde maken een verhoging met meerdere instructies atomair. Stel dat twee threads een gedeelde teller verhogen die aanvankelijk vijf is. Een mogelijke volgorde is: A leest vijf, B leest vijf, A schrijft zes, B schrijft zes. Elke lees- en schrijfbewerking kan coherent zijn, maar de eindstand is zes in plaats van zeven. Het programma heeft een atomaire update of geschikte synchronisatie nodig om beide verhogingen te behouden.

Een cachecontroller die meeluistert, houdt coherentieverzoeken op een verbinding die berichten aan iedereen verspreidt in de gaten. Om een protocol te doorlopen houd je voor de kopie van een blok in elke processorkern een toestand bij en werk je na elke toegang alle betrokken kopieën bij.

Lees de toestanden als toegangsrechten

In het basisprotocol MSI verstuurt een leesmisser een leesverzoek over de bus en laadt het blok in S. Een blok in S schrijven vereist een opwaardering om andere kopieën ongeldig te maken. Schrijven vanuit I vereist een exclusief leesverzoek om gegevens op te halen en eigenaarschap te verkrijgen. De schrijver eindigt in M. Een treffer in M vereist geen coherentietransactie.

Volg vijf toegangen

Begin met X = 0 in het geheugen en zowel A als B in Invalid.

1. A leest X. Een leesverzoek op de bus levert nul. Toestanden: A = S, B = I.
2. B leest X. Een volgend leesverzoek op de bus levert nul. Toestanden: A = S, B = S.
3. A schrijft één. De opwaardering maakt B ongeldig. Toestanden: A = M, B = I.
4. B leest X. B krijgt een misser; A levert één en gaat terug naar S. In het basisprotocol MSI wordt ook het geheugen actueel. Toestanden: A = S, B = S.
5. B schrijft twee. De opwaardering maakt A ongeldig. Toestanden: A = I, B = M.

Elke toegang in dit verloop vereist een bustransactie. De vierde toegang is de belangrijke controle van eigenaarschap: het geheugen alleen bevatte nog de oude waarde voordat A zijn gewijzigde gegevens leverde.

Wat Exclusive bespaart

MESI voegt E toe voor één schone kopie in een cache. Als A een blok leest dat niet in een cache staat en geen andere cache het heeft, krijgt A E. Een volgende schrijfbewerking verandert E zonder bericht in M, omdat er geen andere kopieën ongeldig hoeven te worden gemaakt. MSI zou eerst S laden en daarna voor die schrijfbewerking een opwaardering versturen.

Als een andere processorkern leest terwijl A E heeft, gaan beide naar S. M en E hebben beide één geldige eigenaar, maar alleen M bevat gegevens die nieuwer zijn dan die in het geheugen.

Herhaald schrijven en schaal

Een schrijver die M heeft verkregen, kan het blok tien keer bijwerken zonder tien nieuwe ongeldigmakingen. Een systeem met schrijfbijwerking zou de gewijzigde waarden naar de caches met een gedeelde kopie sturen.

Meeluisteren op berichten aan iedereen wordt kostbaar naarmate meer processorkernen elk verzoek volgen. Coherentieregisters verminderen die verspreiding. Uitgebreidere protocollen voegen nuttige toestanden toe: de toestand Owned van MOESI staat een gewijzigde gedeelde kopie toe en MESIF wijst een cache met de toestand Forward aan om verzoeken te beantwoorden.

Elk coherentieverzoek naar iedereen versturen wordt op grote machines kostbaar. Een coherentieregister houdt bij welke caches elk blok mogelijk bevatten, zodat verzoeken naar die caches kunnen gaan. Hetzelfde bijhouden op blokniveau verklaart ook een veelvoorkomend prestatieprobleem: schijndeling.

Vraag het beheerknooppunt van een blok

Elk blok heeft een item in het coherentieregister van zijn beheerknooppunt dat vastlegt of het niet in een cache, gedeeld of gewijzigd is en welke caches het bevatten. Beheerknooppunten zijn verdeeld over knooppunten of delen van coherentieregisters, waardoor opslag en afhandeling van verzoeken worden gespreid.

Een leesmisser gaat naar het beheerknooppunt. Als het geheugen de actuele gegevens heeft, kan het die leveren en registreert het coherentieregister de nieuwe cache met een gedeelde kopie. Als een andere cache eigenaar is van een kopie in Modified, stuurt het beheerknooppunt het verzoek door naar die eigenaar zodat actuele gegevens kunnen worden geleverd. In een basisprotocol met een coherentieregister wordt hierdoor ook het geheugen weer actueel.

Om te schrijven zonder al eigenaar in Modified te zijn, vraagt een processorkern het beheerknooppunt om toestemming. Het coherentieregister stuurt ongeldigmakingen naar de geregistreerde caches met een gedeelde kopie, wacht op bevestigingen en kent exclusief eigenaarschap toe. Herhaalde schrijftreffers in M vereisen daarna geen verdere coherentieberichten. De tussenkomst van het coherentieregister voegt latentie toe, maar verkeer tussen afzonderlijke punten voorkomt dat elke cache over elk blok wordt bevraagd. Filters voor meeluisterverzoeken op de chip dienen een verwant doel door verzoeken te beperken tot mogelijke houders.

Schat de opslag voor coherentieregisters

Een volledige aanwezigheidsvector heeft één bit per processorkern per blok. Voor 64 kernen is dat 64 bits, oftewel acht bytes. Ten opzichte van een gegevensblok van 64 bytes kosten alleen de aanwezigheidsbits $8\mathrm{/}64=\mathrm{0,125}$, oftewel 12,5%. Bij 256 kernen kost de vector van 32 bytes 50%. Toestandsbits komen daar nog bij. Beperkte lijsten van caches met een gedeelde kopie, grovere vectoren en alleen blokken in caches bijhouden verminderen deze opslaglast.

Herken schijndeling

Werkelijke deling treedt op als processorkernen gegevens via hetzelfde woord uitwisselen. Schijndeling treedt op als ze verschillende variabelen binnen één cacheblok gebruiken en minstens één kern schrijft. De variabelen zijn onafhankelijk in het programma, maar de coherentie-eenheid is groter dan elk van de variabelen.

Twee threads verhogen afzonderlijke tellers van 8 bytes op 0x1000 en 0x1008. Bij blokken van 64 bytes zijn hun bloknummers $\lfloor 4096\mathrm{/}64\rfloor =64$ en $\lfloor 4104\mathrm{/}64\rfloor =64$. Elke verandering van eigenaarschap verplaatst het toegangsrecht voor beide tellers samen. Veelvuldig afwisselend schrijven kan het blok tussen caches heen en weer laten gaan, hoewel geen van beide threads de waarde van de andere nodig heeft.

Als je de tweede teller naar 0x1040 verplaatst, komt die in blok 65. Opvulling of uitlijning kan gegevens per thread scheiden; waarden groeperen per schrijver helpt ook. Een lokale accumulator die één keer wordt teruggeschreven vermindert eigendomsverkeer zonder aan elke waarde opvulling toe te voegen.

Een schrijver en een lezer van een naastgelegen variabele kunnen ook last hebben van schijndeling: de schrijver maakt het verder bruikbare blok van de lezer ongeldig. Twee lezers alleen veroorzaken dit niet. Schijndeling voegt verkeer en vertraging toe; het zorgt er op zichzelf niet voor dat updates van afzonderlijke variabelen onjuiste waarden opleveren. Meet toegangspatronen en plaatsing van cachelijnen voordat je een gegevensindeling verandert.

Een cluster verbindt complete computers, elk met een eigen geheugen en besturingssysteem. Knooppunten werken samen via expliciete berichten. Gegevensplaatsing en communicatiekosten worden daardoor onderdeel van het prestatiemodel van het programma.

Verdeel gegevens en combineer resultaten

MPI, de Message Passing Interface, geeft elk proces een procesnummer en operaties zoals verzenden, ontvangen, verspreiden en reduceren. Een SPMD-programma kan procesnummers gebruiken om de gegevens van elk proces te selecteren. Een pointer in de eigen adresruimte van één knooppunt is geen gewone verwijzing via gedeeld geheugen naar gegevens van een ander knooppunt.

Als één miljoen waarden al over 100 knooppunten zijn verdeeld, telt elk knooppunt lokaal 10.000 waarden op. Een boomreductie combineert de deelsommen in zeven rondes, omdat $\lceil {\log }_{2}100\rceil =7$. Die rondes omvatten netwerkberichten, niet alleen lokale optellingen. Door gegevens dicht bij het werk te houden verminder je overdrachten.

Clusters schalen met standaardcomponenten en defecte hardware kan worden vervangen. Een toepassing kan na het uitvallen van een knooppunt nog steeds herstel of een herstart nodig hebben; vervangbare knooppunten alleen bewaren de voortgang van een onderbroken taak niet.

Meet meer dan de verbindingssnelheid

Latentie is de tijd die een bericht nodig heeft om aan te komen. Verbindingsbandbreedte is de gegevenssnelheid van één verbinding. Totale netwerkbandbreedte telt verbindingscapaciteit op en beschrijft een ideaal totaal waarbij alle verbindingen bezet zijn. Bisectiebandbreedte telt de capaciteit over de doorsnede met de laagste capaciteit die knooppunten in gelijke helften verdeelt.

Vergelijk voor verbindingen met gelijke bandbreedte $B$ netwerken met 64 knooppunten volgens één consistente conventie voor het tellen van verbindingen:

Het rooster heeft acht rijen van zeven horizontale verbindingen en acht kolommen van zeven verticale verbindingen, samen 112. Door beide dimensies met tegenoverliggende randen te verbinden ontstaan de 128 verbindingen van het torusnetwerk. Elk knooppunt van de hyperkubus heeft zes buren, wat $64(6)\mathrm{/}2=192$ verbindingen geeft omdat elke verbinding bij twee eindpunten wordt geteld.

Volg de doorsnede en bekijk daarna de kosten

Een ring die in twee helften wordt verdeeld, kruist slechts twee verbindingen. Veel berichten kunnen daardoor bij die grens concurreren ondanks de in totaal 64 verbindingen van de ring. Een volledig verbonden netwerk kruist $32(32)=1024$ verbindingen tussen zijn twee helften, maar heeft in totaal $64(63)\mathrm{/}2=2016$ verbindingen nodig.

Meer verbindingen kunnen de bandbreedte verbeteren en het aantal tussenstappen verminderen, maar ook het aantal schakelaarpoorten en kabels en het vermogen vergroten. Gedeelde bussen hebben één gedeeld communicatiemiddel; roosters en torusnetwerken spreiden verbindingen over de ruimte. Grote clusters gebruiken vaak geschakelde netwerken met meerdere lagen, zoals verbrede bomen. Hun nut hangt af van het verkeerspatroon en van hun totale capaciteit. Voor communicatie tussen helften van een machine laat de bisectie een grens zien die het totaal verbergt.

Een computer op datacenterschaal combineert tienduizenden servers met stroomvoorziening, koeling en een netwerk om internetdiensten te leveren. Onafhankelijke verzoeken bieden parallel werk, terwijl raamwerken voor batchverwerking zoals MapReduce grote gegevensverwerkingstaken over machines verdelen.

Stem het ontwerp af op de dienst

Systemen op datacenterschaal leggen de nadruk op kosten per verzoek, energie-efficiëntie en beschikbaarheid. Software probeert werk opnieuw, repliceert gegevens en leidt verkeer om defecte machines heen. Een supercomputer legt vaker de nadruk op één nauw gekoppeld wetenschappelijk programma en een snelle verbinding.

Servers zijn verbonden met schakelaars boven in de racks, die via hogere schakelaarniveaus zijn verbonden. Extra niveaus verhogen de communicatielatentie. Verbindingen naar boven zijn vaak overboekt: hun capaciteit ligt onder de gezamenlijke capaciteit van de serververbindingen. Sommige niet-blokkerende netwerken in Clos-stijl bieden voldoende capaciteit op hogere niveaus om deze beperking te vermijden.

Veertig servers met verbindingen van 10 Gbit/s bieden samen 400 Gbit/s aan een rackschakelaar. Vier verbindingen naar boven van 25 Gbit/s leveren 100 Gbit/s, dus de verhouding is 4:1. Als alle servers tegelijk buiten het rack verzenden en gelijk delen, krijgt elke server gemiddeld 2,5 Gbit/s. Een licht belaste server kan meer gebruiken als andere niets doen.

Onderscheid overhead van de faciliteit van nuttig rekenwerk

De effectiviteit van het energiegebruik, PUE, is de totale energie van de faciliteit gedeeld door de energie van de IT-apparatuur over dezelfde periode. Tien MW aan constante IT-belasting plus 1,5 MW aan overhead voor koeling en stroomverdeling geeft $PUE=\mathrm{11,5}\mathrm{/}10=\mathrm{1,15}$. De ideale waarde is 1,0. Google rapporteerde voor 2025 een jaarlijkse PUE van 1,09 over al zijn datacenters samen.

PUE meet de overhead van de faciliteit, niet het nuttige rekenwerk per joule. Servers die niets doen of inefficiënt werken, kunnen binnen de IT-apparatuur energie verspillen, ook als de faciliteit een uitstekende PUE heeft.

Voor 50.000 servers die gemiddeld 300 W gebruiken, is het IT-vermogen 15 MW. Bij PUE 1,2 is het vermogen van de faciliteit 18 MW. Als dat 8.760 uur wordt volgehouden, kost dat 157.680 MWh; alleen de overhead kost 26.280 MWh. Stroomapparatuur, koeling, elektriciteit, servers en gebouwen dragen allemaal bij aan de eigendomskosten.

Ook bij lage belasting kan het vermogen aanzienlijk zijn

Energieproportionaliteit betekent dat het vermogen evenredig met de belasting zou moeten stijgen. Een server met 150 W in rust en 300 W piekvermogen, waarvan het gebruik lineair stijgt, gebruikt bij 30% belasting $150+\mathrm{0,3}(150)=195$ W. De server gebruikt 65% van het piekvermogen voor 30% van de maximale hoeveelheid werk. Eerdere metingen op datacenterschaal lieten zien dat veel servers bij ongeveer 10% tot 50% benutting werkten, waardoor het vermogen in rust belangrijk is.

Houd rekening met storingen en trage reacties

Als 50.000 servers elk gemiddeld één storing per tien jaar hebben, worden voor de hele verzameling 5.000 storingen per jaar verwacht, ongeveer 14 per dag. Die verwachting is een reden voor redundantie en nieuwe pogingen; ze betekent niet dat er elke dag precies 14 storingen zijn.

Een verzoek dat antwoorden van veel servers nodig heeft, wacht op het langzaamste vereiste antwoord. Als elk van 100 servers onafhankelijk 1% kans heeft om traag te zijn, is de kans dat geen ervan traag is ${\mathrm{0,99}}^{100}$. De kans dat minstens één server traag is, is dus $1-{\mathrm{0,99}}^{100}\approx \mathrm{0,634}$. Staartlatentie bij één server wordt bij grote verzoekuitwaaiering een veelvoorkomende vertraging van het hele verzoek.