What happens when a GPU reads memory

Our previous post followed a vector-add kernel — c[i] = a[i] + b[i], 1 thread per float — from nvcc down to the warps. We went into a batch of item connected really the kernel was launched, but we besides near a batch out.

This time, we’re going to reside our omissions, and travel the way the critical SASS instruction (a world load) takes done the hardware — successful this case, since it’s nether my desk, an RTX 4090We do this benignant of reverse engineering for capacity reasons, astatine least in rule (for a awesome rationale, spot 'Why these specifications matter' successful the Citadel microbenchmarking paper). For the same activity applied to much production-relevant GPUs, watch this space.. Little of the item of this way is documented by NVIDIA, astatine slightest not to the level that we’d like, truthful we’ll find it by moving timing experiments connected the hardware itself.

The CUDA kernel we are investigating has 2 lines successful its usability body:

If you inspect the compiled SASS, you’ll spot the instructions that power those lines:

They service to load the elements of the vector bThe instructions are the aforesaid for a, we're pursuing b. from world representation into a register, wherever they tin beryllium added to the elements of a to execute the kernel. One LDG.E asks for 4 bytes successful each of 32 lanes. Serving it takes four 32-byte sectors, 1 cache line, 1 reside translation, a crossbar crossing, one of thirty-six L2 slices, and, erstwhile it misses everywhere, an activate and four file sounds astatine a DRAM chip. It’s this travel of the instruction through the hardware, and back, that we’ll effort to follow.

To group the scene: our warp lives connected 1 of the SM’s 4 sub-partitions, alongside eleven other resident warps. Each rhythm the sub-partition’s scheduler picks 1 warp that is eligible, and issues its adjacent instruction crossed the 32 lanes astatine once. Our warp wins twice: erstwhile for the IMAD.WIDE, and a fewer cycles later (the addresses now sitting successful R4 and R5) for the LDG.

Our communicative starts pinch the LDG.

From the warp to the L1 cache

Let’s commencement pinch the instruction. LDG.E R4, [R4.64] is simply a world load of 32 bits from the 64-bit address stored successful registers R4 and R5R5 appears because of the .64 annotation: registers are 32 bits in size., storing the consequence successful registry R4. To load the information itself, we first must spell get that address from those registers.

One statement of the registry record holds R4 for each 32 lanes astatine onceThe sounds are staged successful an operand collector first. The staging is there for instructions whose sources stock a slope of the registry file, since a bank serves 1 publication per cycle. There are 2 banks, picked by the debased spot of the registry number, truthful an adjacent brace ever spans both.. Another holds R5. The warp sounds some entries, yielding 256 bytes publication arsenic 32 distinct 64-bit addresses, 1 reside per lane.

What the registry retrieve costs

The reside publication adds astatine astir 1 cycle. A shared-memory load taking its address from a registry takes 24 cycles from rumor to first use, and the same load pinch the reside arsenic an contiguous takes 23. (LDG can’t return an immediate).

With each of its addresses resolved, the instruction issues to the load/store portion (LSU). The LSU takes the instruction and its operand addresses, does immoderate reside arithmetic (if necessary)This portion tin adhd contiguous offsets ([R4.64] carries nary offset to add), and scope loads (LDG names the world model directly)., and sends connected the opcode (‘load these addresses’, successful binary), a 32-bit disguise of progressive lanes, its computed addresses, and the number of the registry the consequence belongs in. The adjacent destination is the coalescer.

Each LDG.E instruction successful each lane asks for 4 bytes, but our next destination, the L1 cache, is addressed successful 32 byte sectors. The coalescer’s occupation is to fig retired the minimal number of L1 sectors it needs to retrieve to work our 4-byte requests.

The coalescer figures retired that it ought to emit 4 contiguous assemblage requests, for the 128 bytes the warp has asked for1.

Entering the L1 cache

The petition for 4 contiguous 32-byte sectors is sent onto the L1 cache.

The L1 cache’s portion of statement is still little granular: 128 byte lines. Our 4 contiguous sectors correspond the 4 parts of a azygous line, truthful a request gets made to L1 for that cache line.

First, we person to find whether that statement is already successful the cache. The cache is divided into groups of slots called setsIn method terms, the L1 cache connected the 4090 is 4 measurement set-associative. Caches dishonesty connected a continuum betwixt afloat associative (any cache statement tin be stored anyplace successful the cache), and 'direct-mapped' (each cache statement tin be stored successful only 1 place)., and a line’s address determines which group it belongs to. A group connected this paper holds 4 slots2, and each carries a tag identifying the statement successful it. The lookup compares all 4 against the tag of the statement it wants. The reside it uses is the virtualPresumably truthful that we don't person to salary translator costs to deed L1. reside utilized successful the program3. The group successful which a statement lands is generated from the line’s virtual reside by a hashing schemeIt's a analyzable parity strategy (see the appendix), not conscionable immoderate portion of the bits, truthful that powerfulness of 2 strided accesses (think columns of a matrix, tensor etc.) don't support hitting the aforesaid sets and churn., which you can reverse engineer4.

If 1 of the 4 tags matches and the sectors we want are successful that slot, the data is publication retired and the load is done5. Because we’re loading each of our data for the first time, our petition misses, and must descend further into the memory system.

How overmuch does an L1 deed cost

An L1 deed returns successful astir 15.4 ns — 40 cycles. The number comes from one thread chasing a limited concatenation done a random permutation of L1-resident lines, pinch the latency chase.

Looking for L2: translation

Virtual memory puts 1 level of indirection betwixt the addresses a program names and the addresses astatine which the hardware stores data. The program gets a contiguous abstraction of its own, and the hardware lays that space out crossed beingness pages nevertheless it likes. Translation is the representation between them.

The L1 we conscionable said to was virtually addressed, truthful we didn’t request to concern ourselves pinch translation. Past this point, we person to commencement speaking the hardware’s connection — an L1 miss has to beryllium translated earlier it leaves the SM6.

The existent mapping betwixt beingness and virtual addresses is established at allocation successful the driver: erstwhile b was allocated, the driver picked physical (2MiB) pages for it and wrote page tables into VRAM signaling the assignment7.

The translator portion takes successful a virtual reside and returns a beingness address, according to those tables. The SM keeps its sixteen astir caller translations in a TLB, shared crossed warps8. The very first load will miss successful this TLB.

What translator costs

We can’t spot immoderate costs to hitting the TLB successful immoderate of the probes we have. Misses cost astir 4.4 ns — eleven cycles. The aforesaid refill costs holds wrong 0.1 ns crossed each the pages this spot tin map, and from immoderate SM, truthful the adjacent level of the translator cache is universal, and very cheap.

Once translator has been performed, what leaves is 1 petition per 128-byte line: now pinch the line’s physical address, on pinch a disguise of the sectors we want from it. Ours is simply a azygous petition pinch each 4 sectors marked9.

The petition proceeds retired of the SM, crossed the crossbar to the L2 cache.

Lost successful L2

The petition runs crossed the crossbar to 1 of 36 2 MiB L2 slices, picked by a somewhat analyzable usability of its beingness address10. Any SM tin deed any slice. All slices tin service successful parallel, truthful the aggregate bandwidth is 36x that of a azygous slice.

Inside a slice, the building is of the aforesaid benignant arsenic the L1. Each portion holds 1024 sets. The group to which a statement belongs is picked by a hash of the line’s physical address. Each group now contains 16 slots: the slices are individually 16 measurement set-associative11. The lines are 128 bytes successful size, the aforesaid as in L1.

The statement is not coming successful L2, since we’ve not fetched it beforeThis is possibly creator license: loading the b vector crossed from host memory complete PCIe mightiness person cached it successful L2. But past we couldn't proceed down to DRAM!. Each slice falls done to 1 of 12 representation controllers — 3 slices per controller. Each representation controller’s occupation is to speak to a single GDDR6X DRAM chip12. Our request gets handed complete to that controller.

What does this cost

An L2 deed costs astir 127 ns — immoderate 330 cycles. Each SM tin manus the crossbar up to two line-requests per cycle, and the 36 slices service independently. The exit-port counter is l1tex__m_l1tex2xbar_req_cycles_active.

Found successful DRAM

The representation controller’s occupation is to load the information from its 2 GiB DRAM chip. It does truthful by issuing commands to DRAM complete a bus.

The DRAM is divided into 2 abstracted buses the controller drives independently, called channels. On each transmission beryllium 16 banks: two-dimensional arrays of representation cells. A slope consists of 65,536 rows. The hardware tin unfastened one row astatine a clip (an activate, expensive), and past return immoderate 32-byte columns from that statement (a read, inexpensive while the statement is open).

The reside is taken isolated 1 past time, to lucifer this representation structure. It picks retired a channel, a bank, a row, and a column. Our 4 sectors are four columns of 1 row13.

So, to service our load, the representation controller must first nonstop 1 activate, and past 4 reads14.

What does a DRAM spot do successful consequence to those commands?

Each DRAM compartment is 1 capacitor down 1 transistor. The transistors of a row stock a wordline, attached to their gates. Each transistor sits between its capacitor and a bitline, which runs on a column, providing a path from each compartment (shared pinch the cells of different rows) to the sense amplifiers. Bits are stored successful the complaint authorities of the capacitor. The capacitors perpetually leak charge, truthful the spot has to region each slope now and then to apical them up.

The building of DRAM. Click a statement to enactment arsenic the statement decoder, releasing charge from the capacitors onto the bitline and into the statement buffer.

bitlinerow decoderwordline

The activate bid triggers the row decoder to thrust that row’s wordline, opening the row’s transistors and driving the complaint from the capacitors in that statement (and only that row) done the bitline into the consciousness amplifiers, which amplify that complaint into full-rail bits and clasp them for the controller to read.

When the publication is issued, its file reside picks retired 256 of these row bits. Reading from the consciousness amplifiers gives america very galore bits astatine once, but we need to serialize them onto the pins that thrust information backmost crossed the bus. There are 16 information pins per channel. The 256 bits of our publication time off connected these pins as PAM4GDDR6X is the GDDR6 standard, pinch this PAM4 signalling added. symbols: each awesome is 1 of 4 voltage levels, carrying two bits, truthful 256 bits complete 16 pins is 16 bits per pin — 8 symbols. The timepiece is sent on a shared ligament truthful that the controller tin sample astatine the correct edges.

The measurement back

These PAM4 bursts are deserialized successful the representation controller, and written into the L2 slice’s line. The results tally backmost done the crossbar, backmost to their SM, and capable their L1 slot. They rendezvous pinch the grounds near by their leaving, and their bytes are written into registry R4 crossed each the lanes.

When the load was issued, a dependency obstruction was set, which this register write clears. The warp becomes eligible again, and connected the scheduler’s adjacent cycle it wins the arbitration. The instruction it issues is the adhd that was waiting on b[i].

The information travel — L1, TLB, crossbar, L2, controller, and backmost — costs astir 255 ns, immoderate 660 cycles. All the while our warp was parked connected its barrier. The remainder of the spot wasn’t idle though. The sub-partition issued the same loads for different 11 warps, the remainder of the SM for different 36, the other SMs for the different 6096. The consequence is simply a cacophony of loads, the per-load latency of immoderate 1 of them mislaid successful the noise. Here’s what that looks like:

A timing-proportional simulation of the execution of only the instructions in the vadd kernel that correspond to the load of b. Each SM loads only those addresses it loads successful the existent kernel: those addresses ray up (and miss) successful the correct L1 set, past are routed done the crossbar to the correct L2 slice, wherever they miss, falling done a correctly contended representation controller to a simulated DRAM bank, before returning backmost done L2, backmost done L1, and returning their results into the correct register.

Appendix: the probes

Setup

All measurements are connected 1 RTX 4090 (sm_89), pinch the halfway timepiece locked at 2.6 GHz. Cycles travel from measured nanoseconds astatine that frequency. Two main instruments:

A latency chase. To get a latency measurement (especially erstwhile that latency changing tells you thing astir the chip), we tally a pointer cycle through a chosen group of lines, hopped 20,000 times, and past measurement the mean ns per hop. If the lines we constituent to fresh successful a cache level, then they enactment resident, and the mean is that level’s deed latency. Because of the steepness of the hierarchy, immoderate loads that overflow to the adjacent level down tend to show up powerfully successful the average. ld.global.ca (LDG.E…STRONG.SM) for chases astatine the L1, ld.global.cg (LDG.E…STRONG.GPU) goes past L1. Hit latencies are 15.4 ns astatine the L1, 127.4 ns astatine the L2, and 255.4 ns astatine DRAM.

Hardware counters. To publication ncu’s counters reliably you person to return them arsenic slopes complete loop count truthful fixed overhead cancels. Sector and petition counters astatine the L1 exit larboard and the L2 broadside are utilized to figure retired much astir the style of the requests, and a per-slice assemblage counter helps to springiness america the L2 portion measurements.

The L1 group function

The 8 bits of the L1 scale are the XOR of a fixed subset of the reside bits. Written as a bitmask complete the address, 1 ground for those subsets is:

bitmaskbitmask

0	0xc3901e00		4	0x47810400
1	0x119a80a00		5	0x1b4e09180
2	0x167041b00		6	0xb6405400
3	0xdbc21d80		7	0xdc202c80

The masks themselves aren’t unsocial — immoderate invertible operation of these 8 describes the same partition.

Page tables and the TLB

The 16-entry TLB is only the first level, but what happens erstwhile you miss? A miss refills successful astir 4.4 ns, and an L2 deed is 127 ns and a VRAM entree is 255 ns, truthful we can’t beryllium going from those. The conclusion is that it comes from some larger on-chip translator cache.

The costs is level wrong 0.1 ns for each the pages the spot tin map, and from any SM. More evidence: stepping the page tables with nvdebug shows the volatile spot group connected each directory entry, truthful they’re not cached successful the normal hierarchy.

The L2 portion function

Measuring which portion owns a statement is beautiful hard. The L2 is physically indexed, truthful the probe has to activity successful device-physical addresses from the page-table walk. Nsight Compute does person a per-slice assemblage counter, but reports only the min, max, average, and sum crossed the 36 instances, ne'er the existent portion index.

Even so, the aggregate is capable to show whether 2 addresses stock a slice. If the 2 addresses unrecorded connected the aforesaid slice, aft loading both, the max counter reports 2, if they’re connected different slices the max is 1. You tin use this probe to get a typical reside that lands connected each of the 36 slices.

With the 36 representatives successful hand, you tin get immoderate caller candidate’s slice. If you publication the campaigner galore times alongside each 36, pinch each of the different addresses publication a chopped number of times (say 20001, 20002, … times), the sum of the candidate’s publication count and only 1 of the representatives will lucifer the max counter, and you tin fig the portion by inference.

From that, you tin nutrient a array of galore beingness address-slice pairs. The hard portion is going from specified a array to a physically plausible function. One tool that helped america a spot was moving the aforesaid kinds of experiments connected two different chips built connected the aforesaid die: the 4090, and the L40S, which has an extra slice per representation controller.

Here’s 1 Claude made earlierIt's difficult to beryllium judge what's really successful the hardware here, but this is plausible fixed my constricted knowledge. The priors: there's sewage to beryllium some shared silicon betwixt the L40S and the 4090 (assuming NVIDIA don't vessel two completely different functional paths for chips connected the aforesaid layout but with different amounts of L2 fused off). And the usability has to beryllium simple-ish in hardware, i.e. XORs, arithmetic etc. are adjacent game, but if Claude tries to put in a 4096 introduction lookup array you show it to spell effort harder.:

Whilst it is very difficult to find specified a function, it’s very easy to show if you’ve recovered 1 that works. Drawing 8,192 L2-resident lines from precisely k predicted slices:

lines drawn fromMload/svs k=1

1 predicted slice	1,957	1.00×
2	3,917	2.00×
4	7,826	4.00×
9	17,582	8.98×
18	34,446	17.60×
all 36	68,085	34.78×

The L2 group scale and geometry

Once the portion usability pins addresses to a azygous slice, you tin do the same eviction-set archaeology connected that slice, to fig retired the structure, which tells you that it’s 16 measurement set-associative (a pursuit pinch 17 elements thrashes, but one pinch 16 doesn’t).

The group scale wrong a portion is the aforesaid benignant of parity usability arsenic the L1’s — ten bits, pinch the aforesaid (a >> 15) mod 9 nonlinearity successful the apical bit. Unfortunately, the masks progressive disagree depending connected the slice. For 1 slice:

It has immoderate properties that fto you sense-check it. For example: a contiguous 72MiB fills each slot successful each portion without thrashing anything, arsenic you’d expect.

DRAM refresh

DRAM cells leak complaint and truthful person to beryllium periodically refreshed, which makes some kinds of timing probes harder. You tin spot it by moving a limited chase that writes each hop’s timing into shared memory. Most DRAM accesses travel back at the accustomed latency, but a mini stock return longer, dispersed evenly retired to a hard ceiling astir 210 ns higher than usual. An evenly spaced tally for illustration that is the signature of a fixed magnitude stall. The stall is ~210 ns. About 2% of accesses deed one. It doesn’t deed the full spot astatine erstwhile — it’s much section than that — but I couldn’t show what the portion was.