Wat er gebeurt wanneer een GPU geheugen leest

Ditmaal vullen we die gaten in en volgen we het pad dat een kritieke SASS-instructie (een global load) aflegt door de hardware — in dit geval een RTX 4090. We doen dit soort reverse engineering in principe om prestatieredenen (voor een goede onderbouwing, zie 'Why these details matter' in het Citadel microbenchmarking-paper).

Omdat NVIDIA weinig details over dit pad documenteert op het niveau dat we wensen, bepalen we dit door timing-experimenten uit te voeren op de hardware zelf.

De CUDA-kernel die we onderzoeken heeft twee regels in de functiebody:

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

Als je de gecompileerde SASS inspecteert, zie je de instructies die deze regels aansturen:

/*0080*/  IMAD.WIDE R4, R6, R7, c[0x0][0x168] ;   // &b[i]
/*00a0*/  LDG.E R4, [R4.64] ;                     // b[i]

Deze dienen om de elementen van vector b vanuit het globale geheugen naar een register te laden, waar ze kunnen worden opgeteld bij de elementen van a om de kernel uit te voeren.

Eén LDG.E vraagt om vier bytes in elk van de 32 lanes. Het verwerken hiervan vereist vier 32-byte sectoren, één cachelijn, één adresvertaling, een crossbar-oversteek, één van de zesendertig L2-slices en, wanneer er overal een 'miss' optreedt, een activatie en vier kolom-uitlezingen bij een DRAM-chip. We zullen deze reis van de instructie door de hardware, en terug, volgen.

De context: onze warp bevindt zich op één van de vier sub-partities van de SM, samen met elf andere aanwezige warps. Elke cyclus kiest de scheduler van de sub-partitie één warp die in aanmerking komt en voert de volgende instructie over de 32 lanes tegelijk uit. Onze warp wint tweemaal: eerst voor de IMAD.WIDE, en een paar cycli later (wanneer de adressen in R4 en R5 staan) voor de LDG.

Ons verhaal begint bij de LDG.

Van de warp naar de L1-cache

De instructie LDG.E R4, [R4.64] is een globale load van 32 bits vanaf het 64-bit adres dat is opgeslagen in registers R4 en R5 (R5 verschijnt vanwege de .64 annotatie: registers zijn 32 bits groot), waarbij het resultaat in register R4 wordt opgeslagen. Om de data zelf te laden, moeten we eerst dat adres uit die registers ophalen.

Eén rij van het registerbestand bevat R4 voor alle 32 lanes tegelijk. De uitlezingen worden eerst in een operand collector geplaatst. Deze staging is bedoeld voor instructies waarvan de bronnen een bank van het registerbestand delen, aangezien een bank één uitlezing per cyclus bedient. Er zijn twee banken, gekozen op basis van het laagste bit van het registernummer, waardoor een aangrenzend paar altijd beide banken beslaat. Een andere rij bevat R5. De warp leest beide vermeldingen, wat resulteert in 256 bytes die worden gelezen als 32 afzonderlijke 64-bit adressen, één adres per lane.

De kosten van de register-uitlezing

De adresuitlezing kost maximaal één cyclus. Een shared-memory load die zijn adres uit een register haalt, doet er 24 cycli over van issue tot eerste gebruik; dezelfde load met het adres als een immediate kost 23 cycli (LDG kan geen immediate gebruiken).

Zodra alle adressen zijn opgelost, wordt de instructie uitgebracht naar de load/store unit (LSU). De LSU neemt de instructie en de operandadressen, voert eventuele adresberekeningen uit (deze unit kan immediate offsets optellen, hoewel [R4.64] geen offset heeft) en definieert scope loads (LDG benoemt direct het globale window). Vervolgens stuurt de LSU de opcode ('laad deze adressen', in binair), een 32-bit masker van actieve lanes, de berekende adressen en het nummer van het register waarin het resultaat moet komen, door naar de coalescer.

Elke LDG.E instructie in elke lane vraagt om 4 bytes, maar onze volgende bestemming, de L1-cache, wordt geadresseerd in sectoren van 32 bytes. De taak van de coalescer is om te bepalen wat het minimale aantal L1-sectoren is dat moet worden opgehaald om aan onze 4-byte verzoeken te voldoen. In dit geval stelt de coalescer vast dat er 4 aaneengesloten sectorverzoeken moeten worden verzonden voor de 128 bytes waar de warp om heeft gevraagd.

De L1-cache binnengaan

Het verzoek om vier aaneengesloten 32-byte sectoren wordt naar de L1-cache gestuurd. De organisatie-eenheid van de L1-cache is nog minder granulair: cachelijnen van 128 bytes. Onze 4 aaneengesloten sectoren vertegenwoordigen de 4 delen van één enkele lijn, dus er wordt een verzoek aan de L1 gedaan voor die specifieke cachelijn.

Eerst moet worden bepaald of die lijn al in de cache aanwezig is. De cache is verdeeld in groepen slots, genaamd sets. Op de RTX 4090 is de L1-cache 4-way set-associative. Caches bevinden zich op een continuüm tussen fully associative (elke cachelijn kan overal in de cache worden opgeslagen) en direct-mapped (elke cachelijn kan slechts op één plek worden opgeslagen).

Het adres van een lijn bepaalt bij welke set hij hoort. Een set op deze kaart bevat vier slots, en elk slot heeft een tag die de lijn identificeert. De lookup vergelijkt alle vier de tags met de tag van de gewenste lijn. Het gebruikte adres is het virtuele adres zoals gebruikt in het programma. De set waarin een lijn terechtkomt, wordt gegenereerd vanuit het virtuele adres via een hash-schema (een complex parity-schema, zie bijlage), zodat toegangen met een macht-van-2 stapgrootte (zoals kolommen van een matrix of tensor) niet constant dezelfde sets raken en deze voortdurend legen (churn).

Als een van de vier tags overeenkomt en de sectoren die we willen in dat slot zitten, worden de gegevens uitgelezen en is de load voltooid. Omdat we onze gegevens voor het eerst laden, is er sprake van een miss en moet het verzoek dieper in het geheugensysteem afdalen.

Wat kost een L1-hit? Een L1-hit is terug in ongeveer 15,4 ns — 40 cycli. Dit getal komt voort uit één thread die een afhankelijke keten volgt door een willekeurige permutatie van L1-resident lijnen, inclusief de latency chase.

Zoeken naar L2: vertaling

Virtueel geheugen plaatst een niveau van indirectie tussen de adressen die een programma benoemt en de adressen waarop de hardware data opslaat. Het programma krijgt een eigen aaneengesloten ruimte, en de hardware verdeelt die ruimte over fysieke pagina's naar eigen inzicht. Vertaling (translation) is de kaart tussen deze twee.

De L1-cache waar we zojuist over spraken was virtueel geadresseerd, dus hoefden we ons geen zorgen te maken over vertaling. Vanaf dit punt moeten we de taal van de hardware spreken — een L1-miss moet worden vertaald voordat hij de SM verlaat.

De daadwerkelijke mapping tussen fysieke en virtuele adressen wordt bij allocatie in de driver vastgesteld: wanneer b werd gealloceerd, koos de driver fysieke pagina's (2MiB) hiervoor en schreef paginatabellen in het VRAM om deze toewijzing vast te leggen. De vertaalunit neemt een virtueel adres en geeft een fysiek adres terug, conform deze tabellen. De SM bewaart de zestien meest recente vertalingen in een TLB, die gedeeld wordt over de warps. De allereerste load zal een miss in deze TLB veroorzaken.

Wat kost vertaling? In de beschikbare probes is er geen merkbare kost voor het raken van de TLB. Misses kosten ongeveer 4,4 ns — elf cycli. Deze refill-kost is consistent binnen 0,1 ns voor alle pagina's die deze chip kan mappen en vanuit elke SM, wat suggereert dat het volgende niveau van de vertaalcache universeel en zeer goedkoop is.

Zodra de vertaling is uitgevoerd, vertrekt er één verzoek per 128-byte lijn: nu met het fysieke adres van de lijn, samen met een masker van de sectoren die we daaruit willen. In ons geval is dit één verzoek waarbij alle vier de sectoren zijn gemarkeerd. Het verzoek verlaat de SM en gaat via de crossbar naar de L2-cache.

Verdwaald in L2

Het verzoek gaat via de crossbar naar één van de 36 L2-slices van 2 MiB, gekozen door een complexe functie van het fysieke adres. Elke SM kan elke slice bereiken. Alle slices kunnen parallel werken, waardoor de geaggregeerde bandbreedte 36 keer die van een enkele slice is.

Binnen een slice is de structuur vergelijkbaar met die van de L1. Elke slice bevat 1024 sets. De set waartoe een lijn behoort, wordt gekozen via een hash van het fysieke adres van de lijn. Elke set bevat nu 16 slots: de slices zijn individueel 16-way set-associative. De lijnen zijn 128 bytes groot, net als in de L1.

De lijn is niet aanwezig in de L2, aangezien we deze nog niet eerder hebben opgehaald. Elke slice valt dan terug op één van de 12 geheugencontrollers — 3 slices per controller. De taak van elke geheugencontroller is om te communiceren met één enkele GDDR6X DRAM-chip. Ons verzoek wordt overgedragen aan die controller.

Wat kost dit? Een L2-hit kost ongeveer 127 ns — circa 330 cycli. Elke SM kan de crossbar voorzien van maximaal twee lijnverzoeken per cyclus, en de 36 slices dienen onafhankelijk.

Gevonden in DRAM

De geheugencontroller moet de data laden uit zijn 2 GiB DRAM-chip. Dit gebeurt door commando's naar de DRAM te sturen over een bus.

De DRAM is verdeeld in twee aparte bussen die de controller onafhankelijk aanstuurt, genaamd channels. Op elk kanaal bevinden zich 16 banks: tweedimensionale arrays van geheugencellen. Een bank bestaat uit 65.536 rijen. De hardware kan één rij tegelijk openen (een activate, wat duur is), en kan vervolgens elke 32-byte kolom uit die rij teruggeven (een read, wat goedkoop is zolang de rij open is).

Het adres wordt voor de laatste keer ontleed om aan deze geheugenstructuur te voldoen. Er wordt een kanaal, een bank, een rij en een kolom gekozen. Onze vier sectoren zijn vier kolommen van één rij. Om aan onze load te voldoen, moet de geheugencontroller dus eerst één activate sturen, gevolgd door vier reads.

Hoe de DRAM antwoordt

Wat doet een DRAM-chip in reactie op deze commando's? Elke DRAM-cel is een condensator achter één transistor. De transistors van een rij delen een wordline, verbonden met hun gates. Elke transistor bevindt zich tussen zijn condensator en een bitline, die langs een kolom loopt en een pad biedt van elke cel (gedeeld met cellen van andere rijen) naar de sense amplifiers. Bits worden opgeslagen in de ladingsstatus van de condensator. Omdat condensatoren constant lading lekken, moet de chip af en toe elke bank pauzeren om ze weer op te laden.

Het activate-commando activeert de row decoder om de wordline van die rij aan te sturen. Dit opent de transistors van de rij en drijft de lading uit de condensatoren van die rij (en alleen die rij) via de bitline naar de sense amplifiers. Deze versterken die lading tot volledige bits en houden deze vast zodat de controller ze kan lezen.

Wanneer de read wordt uitgebracht, selecteert het kolomadres 256 van deze rij-bits. Het lezen uit de sense amplifiers geeft zeer veel bits tegelijk, maar deze moeten worden geserialiseerd op de pinnen die de data over de bus terugsturen. Er zijn 16 datapinnen per kanaal. De 256 bits van onze read vertrekken via deze pinnen als PAM4-symbolen: elk symbool is een van vier spanningsniveaus en draagt twee bits. 256 bits over 16 pinnen betekent 16 bits per pin — acht symbolen. De klok wordt via een gedeelde draad verzonden, zodat de controller op de juiste flanken kan samplen.

De weg terug

Deze PAM4-bursts worden gedeserialiseerd in de geheugencontroller en weggeschreven in de lijn van de L2-slice. De resultaten gaan terug via de crossbar naar hun SM en vullen hun L1-slot. Ze komen samen met het record dat is achtergelaten bij hun vertrek, en hun bytes worden geschreven in register R4 over alle lanes.

Toen de load werd uitgebracht, werd een dependency barrier ingesteld, die door deze register-schrijfactie wordt opgeheven. De warp wordt weer in aanmerking genewcommand en bij de volgende cyclus van de scheduler wint hij de arbitrage. De instructie die hij vervolgens uitvoert, is de optelling die wachtte op b[i].

De volledige round trip — L1, TLB, crossbar, L2, controller en terug — kost ongeveer 255 ns, circa 660 cycli. De hele tijd stond onze warp geparkeerd op zijn barrier. De rest van de chip was echter niet inactief. De sub-partitie voerde dezelfde loads uit voor andere 11 warps, de rest van de SM voor nog eens 36, en de andere SM's voor de overige 6096. Het resultaat is een cacofonie van loads, waarbij de latency per individuele load verloren gaat in het lawaai.

***

Bijlage: de probes

Setup

Alle metingen zijn uitgevoerd op één RTX 4090 (sm_89), met de core clock gefixeerd op 2,6 GHz. Cycli zijn afgeleid van gemeten nanoseconden bij die frequentie. Er zijn twee hoofdinstrumenten gebruikt:

  1. Een latency chase: Om een latency-meting te krijgen, draaien we een pointer cycle door een gekozen set lijnen, 20.000 keer gesprongen, en meten we de gemiddelde ns per sprong. Als de lijnen in een bepaald cacheniveau passen, blijven ze resident en is het gemiddelde de hit-latency van dat niveau.
  • ld.global.ca (LDG.E…STRONG.SM) voor chases in de L1.
  • ld.global.cg (LDG.E…STRONG.GPU) gaat voorbij de L1.
  • Hit latencies: 15,4 ns (L1), 127,4 ns (L2), en 255,4 ns (DRAM).
  1. Hardware counters: Om de counters van ncu betrouwbaar uit te lezen, worden ze genomen als hellingen over het aantal iteraties zodat vaste overhead wegvalt.

De L1 set-functie

De 8 bits van de L1-index zijn de XOR van een vaste subset van de adresbits. Een basis voor deze subsets, geschreven als een bitmasker over het adres, is:

IndexBitmaskIndexBitmask
00xc3901e0040x47810400
10x119a80a050x1b4e09180
20x167041b060xb6405400
30xdbc21d8070xdc202c80

Paginatabellen en de TLB

De TLB met 16 vermeldingen is slechts het eerste niveau. Een miss wordt hersteld in ongeveer 4,4 ns. Aangezien een L2-hit 127 ns kost en VRAM-toegang 255 ns, kan de vertaling niet daarvandaan komen. De conclusie is dat dit komt uit een grotere on-chip vertaalcache.

De L2 slice-functie

Het bepalen van welke slice een lijn bezit is zeer complex. De L2 is fysiek geïndexeerd, dus de probe moet werken met fysieke apparaadadressen uit de paginatabel-walk. Met behulp van per-slice sector counters in Nsight Compute kan worden bepaald of twee adressen dezelfde slice delen.

Hieronder volgt een plausibele implementatie van de slice-functie in Python:

SHIFT, OFFSET = (5, 0, 1), (1, 0, 0)

def parity(x):
    return bin(x).count("1") & 1

def _state(a, N):
    wide = (N == 48)                             # L40S: 4 slices/controller
    b35 = (1 << 35) if wide else 0
    # stage 1 — welke van de 12 controllers
    P1c = parity(a & 0x76A990400)                # controller parity 1
    P1  = parity(a & (0x76A990400 ^ b35))        # wide form voor L40S
    P2  = parity(a & 0x2CCF7B000)                # controller parity 2
    A   = ((a >> 15) + 2*parity(a & 0x3C9041000) + parity(a & (0x2882B0800 ^ b35)) + 2) % 3
    # stage 2 — welke slice binnen de controller
    g   = ((a + (1 << 16)) >> 17) % 9            # cyclic counter
    q0  = parity(a & 0x8000)                     # correctie parities
    q1  = parity(a & 0x5985E0500)
    q2  = parity(a & (0x2354E4400 ^ b35))
    q3  = parity(a & 0x3C9041000)
    carry = 1 if q0 + q1 + q2 >= 2 else 0
    start = (5 + 7*q0 + 5*q1 + 2*q2 + q3 - carry) % 9
    o     = (g - SHIFT[A] - start) % 9
    Lf    = 2 if (q0 ^ q1 ^ q2) == 0 else 1
    return P1c, P1, P2, A, q2, o // 3, (1 if (o % 3) >= Lf else 0)

def slice_of(a, N=36):
    P1c, P1, P2, A, q2, d, u = _state(a, N)
    controller = (2*P1c + P2) * 3 + A            # 0..11
    if N == 36:                                  # 4090: 3 slices per controller
        base = 2 if d == 0 else (1 if (d == 1 and u == 0) else 0)
        B = ((1 - base) % 3 if q2 else base) % 3
        B = (B + OFFSET[A]) % 3
        return controller * 3 + B
    if N == 48:                                  # L40S: 4 slices per controller
        i0, i1 = P1 ^ q2 ^ u, P1 ^ P2 ^ u
        return controller * 4 + 2*i0 + i1
    raise ValueError("N must be 36 or 48")

De effectiviteit van deze functie kan worden getoetst door te kijken hoeveel lijnen uit voorspelde slices worden geladen:

Voorspelde slicesGeladen lijnenFactor vs k=1
1 voorspelde slice1.9571,00x
23.9172,00x
47.8264,00x
917.5828,98x
1834.44617,60x
Alle 3668.08534,78x

De L2 set-index en geometrie

Zodra adressen aan een specifieke slice zijn gekoppeld, kan via eviction-set archeologie worden vastgesteld dat de L2 16-way set-associative is. De set-index binnen een slice is een parity-functie vergelijkbaar met die van de L1. Voor één specifieke slice ziet de functie er als volgt uit:

def parity(x):
    return bin(x).count("1") & 1

def set_index(a):
    q = a // 1152
    b0 = parity(a & 0x0bd654c80) ^ parity(q & 0x00e500)
    b1 = parity(a & 0x0bd654c80) ^ parity(q & 0x010000)
    b2 = parity(a & 0x07aed8b80) ^ parity(q & 0x027c00)
    b3 = parity(a & 0x03e313180) ^ parity(q & 0x045500)
    b4 = parity(a & 0x03e313300) ^ parity(q & 0x080300)
    b5 = parity(a & 0x0bd654e80) ^ parity(q & 0x104200)
    b6 = parity(a & 0x0bd654c00) ^ parity(q & 0x200b00)
    b7 = parity(a & 0x044dcb880) ^ parity(q & 0x401600)
    b8 = parity(a & 0x000000200) ^ parity(q & 0x804600)
    b9 = parity(a & 0x13bc21180) ^ parity(q & 0x006400) ^ int((a >> 15) % 9 in (2, 6))
    return sum(b << i for i, b in enumerate([b0, b1, b2, b3, b4, b5, b6, b7, b8, b9]))

DRAM-refresh

DRAM-cellen lekken lading en moeten periodiek worden ververst (refreshed). Dit is zichtbaar in timing-probes: de meeste DRAM-toegangen hebben de gebruikelijke latency, maar een klein deel (circa 2%) duurt langer, met een vaste stall van ongeveer 210 ns.

***

Voetnoten

  1. Sectoren per load: Bij een stride van 1 beslaan 32 lanes 128 aaneengesloten bytes, wat resulteert in vier sectoren.
  2. Slots in een set: Bepaald door een pool van kandidaat-lijnen te maken die groter is dan de L1-cache en te kijken wanneer de latency plotseling daalt bij het verwijderen van adressen.
  3. Virtuele adressering: De L1 gebruikt het virtuele adres voor zowel index als tag. Dit is bewezen door één fysieke allocatie op twee virtuele adressen te mappen.
  4. Index-bits: De maskers zijn bepaald door minimale conflictsets binnen één 2 MiB pagina en door bits boven de pagina te flippen.
  5. Sectoren in slots: Een slot kan een lijn bevatten waarbij slechts enkele sectoren aanwezig zijn. L2-counters bevestigen dat er alleen sectoren worden aangevraagd die de lanes daadwerkelijk raken.
  6. L2 fysieke tagging: Bewezen door te mappen van één fysieke allocatie op twee virtuele adressen; de L2-threshold wordt bij beide op hetzelfde fysieke formaat bereikt.
  7. GPU paginatabellen: Direct leesbaar via tools zoals nvdebug. De meeste allocaties eindigen in een 2 MiB paginatabel-entry.
  8. TLB eigenschappen: De TLB is fully associative, bevat 16 entries per SM en gebruikt een LRU-vervangingsbeleid.
  9. L2 requests: Vergelijkbare logica als bij de L1, maar gebruikmakend van L2-counters.
  10. Slices en controllers: De 4090 heeft 12 controllers met elk 3 actieve slices (één slice is gefused). Dit is afgeleid uit metingen en vergelijkingen met de L40S.
  11. L2 archeologie: Gआधारd op eviction-sets, met de restrictie dat alleen adressen die naar dezelfde slice mappen gebruikt kunnen worden.
  12. Busbreedte: De 12 controllers corresponderen met de 384-bit bus (12 × 32-bit).
  13. Rijgrootte: Bepaald door timing-verschillen tussen toegang tot dezelfde rij versus verschillende rijen. De rijgrootte is 1 KiB (32 kolommen van 32 bytes).
  14. Activatiekosten: Het openen van een nieuwe rij kost ongeveer 15x meer dan een read in een reeds open rij (ongeveer 3,4 ns per extra read).