What happens when a GPU reads memory

Our previous post followed a vector-add kernel — c[i] = a[i] + b[i], one thread per float — from nvcc down to the warps. We went into a lot of detail on how the kernel was launched, but we also left a lot out.

This time, we’re going to address our omissions, and follow the path the critical SASS instruction (a global load) takes through the hardware — in this case, since it’s under my desk, an RTX 4090We do this kind of reverse engineering for performance reasons, at least in principle (for a great rationale, see 'Why these details matter' in the Citadel microbenchmarking paper). For the same work applied to more production-relevant GPUs, watch this space.. Little of the detail of this path is documented by NVIDIA, at least not to the level that we’d like, so we’ll determine it by running timing experiments on the hardware itself.

The CUDA kernel we are investigating has two lines in its function body:

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

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

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

They serve to load the elements of the vector bThe instructions are the same for a, we're following b. from global memory into a register, where they can be added to the elements of a to perform the kernel. One LDG.E asks for four bytes in each of 32 lanes. Serving it takes four 32-byte sectors, one cache line, one address translation, a crossbar crossing, one of thirty-six L2 slices, and, when it misses everywhere, an activate and four column reads at a DRAM chip. It’s this journey of the instruction through the hardware, and back, that we’ll try to follow.

To set the scene: our warp lives on one of the SM’s four sub-partitions, alongside eleven other resident warps. Each cycle the sub-partition’s scheduler picks one warp that is eligible, and issues its next instruction across the 32 lanes at once. Our warp wins twice: once for the IMAD.WIDE, and a few cycles later (the addresses now sitting in R4 and R5) for the LDG.

Our story starts with the LDG.

From the warp to the L1 cache

Let’s start with the instruction. LDG.E R4, [R4.64] is a global load of 32 bits from the 64-bit address stored in registers R4 and R5R5 appears because of the .64 annotation: registers are 32 bits in size., storing the result in register R4. To load the data itself, we first must go get that address from those registers.

One row of the register file holds R4 for all 32 lanes at onceThe reads are staged in an operand collector first. The staging is there for instructions whose sources share a bank of the register file, since a bank serves one read per cycle. There are two banks, picked by the low bit of the register number, so an adjacent pair always spans both.. Another holds R5. The warp reads both entries, yielding 256 bytes read as 32 distinct 64-bit addresses, one address per lane.

With all of its addresses resolved, the instruction issues to the load/store unit (LSU). The LSU takes the instruction and its operand addresses, does some address arithmetic (if necessary)This unit can add immediate offsets ([R4.64] carries no offset to add), and scope loads (LDG names the global window directly)., and sends on the opcode (‘load these addresses’, in binary), a 32-bit mask of active lanes, its computed addresses, and the number of the register the result belongs in. The next destination is the coalescer.

Each LDG.E instruction in each lane asks for 4 bytes, but our next destination, the L1 cache, is addressed in 32 byte sectors. The coalescer’s job is to figure out the minimal number of L1 sectors it needs to retrieve to service our 4-byte requests.

The coalescer figures out that it ought to emit 4 contiguous sector requests, for the 128 bytes the warp has asked for¹.

Entering the L1 cache

The request for four contiguous 32-byte sectors is sent onto the L1 cache.

The L1 cache’s unit of organization is still less granular: 128 byte lines. Our 4 contiguous sectors represent the 4 parts of a single line, so a request gets made to L1 for that cache line.

First, we have to determine whether that line is already in the cache. The cache is divided into groups of slots called setsIn technical terms, the L1 cache on the 4090 is 4 way set-associative. Caches lie on a continuum between fully associative (any cache line can be stored anywhere in the cache), and 'direct-mapped' (each cache line can be stored in only one place)., and a line’s address determines which set it belongs to. A set on this card holds four slots², and each carries a tag identifying the line in it. The lookup compares all four against the tag of the line it wants. The address it uses is the virtualPresumably so that we don't have to pay translation cost to hit L1. address used in the program³. The set in which a line lands is generated from the line’s virtual address by a hashing schemeIt's a complex parity scheme (see the appendix), not just some slice of the bits, so that power of 2 strided accesses (think columns of a matrix, tensor etc.) don't keep hitting the same sets and churn., which you can reverse engineer⁴.

If one of the four tags matches and the sectors we want are in that slot, the data is read out and the load is done⁵. Because we’re loading all of our data for the first time, our request misses, and must descend further into the memory system.

Looking for L2: translation

Virtual memory puts one level of indirection between the addresses a program names and the addresses at which the hardware stores data. The program gets a contiguous space of its own, and the hardware lays that space out across physical pages however it likes. Translation is the map between them.

The L1 we just spoke to was virtually addressed, so we didn’t need to concern ourselves with translation. Past this point, we have to start speaking the hardware’s language — an L1 miss has to be translated before it leaves the SM⁶.

The actual mapping between physical and virtual addresses is established at allocation in the driver: when b was allocated, the driver picked physical (2MiB) pages for it and wrote page tables into VRAM recording the assignment⁷.

The translation unit takes in a virtual address and returns a physical address, according to those tables. The SM keeps its sixteen most recent translations in a TLB, shared across warps⁸. The very first load will miss in this TLB.

Once translation has been performed, what leaves is one request per 128-byte line: now with the line’s physical address, along with a mask of the sectors we want from it. Ours is a single request with all four sectors marked⁹.

The request proceeds out of the SM, across the crossbar to the L2 cache.

Lost in L2

The request runs across the crossbar to one of 36 2 MiB L2 slices, picked by a somewhat complex function of its physical address¹⁰. Any SM can hit any slice. All slices can serve in parallel, so the aggregate bandwidth is 36x that of a single slice.

Inside a slice, the structure is of the same kind as the L1. Each slice holds 1024 sets. The set to which a line belongs is picked by a hash of the line’s physical address. Each set now contains 16 slots: the slices are individually 16 way set-associative¹¹. The lines are 128 bytes in size, the same as in L1.

The line is not present in L2, since we’ve not fetched it beforeThis is perhaps artistic license: loading the b vector across from host memory over PCIe might have cached it in L2. But then we couldn't continue down to DRAM!. Each slice falls through to one of 12 memory controllers — 3 slices per controller. Each memory controller’s job is to speak to a single GDDR6X DRAM chip¹². Our request gets handed over to that controller.

Found in DRAM

The memory controller’s job is to load the data from its 2 GiB DRAM chip. It does so by issuing commands to DRAM over a bus.

The DRAM is divided into two separate buses the controller drives independently, called channels. On each channel sit 16 banks: two-dimensional arrays of memory cells. A bank consists of 65,536 rows. The hardware can open one row at a time (an activate, expensive), and then return any 32-byte columns from that row (a read, cheap while the row is open).

The address is taken apart one last time, to match this memory structure. It picks out a channel, a bank, a row, and a column. Our four sectors are four columns of one row¹³.

So, to serve our load, the memory controller must first send one activate, and then four reads¹⁴.

How the DRAM replies

What does a DRAM chip do in response to those commands?

Each DRAM cell is one capacitor behind one transistor. The transistors of a row share a wordline, attached to their gates. Each transistor sits between its capacitor and a bitline, which runs along a column, providing a path from each cell (shared with the cells of other rows) to the sense amplifiers. Bits are stored in the charge state of the capacitor. The capacitors constantly leak charge, so the chip has to pause each bank now and then to top them up.

The structure of DRAM. Click a row to act as the row decoder, releasing charge from the capacitors onto the bitline and into the row buffer.

The activate command triggers the row decoder to drive that row’s wordline, opening the row’s transistors and driving the charge from the capacitors in that row (and only that row) through the bitline into the sense amplifiers, which amplify that charge into full-rail bits and hold them for the controller to read.

When the read is issued, its column address picks out 256 of these row bits. Reading from the sense amplifiers gives us very many bits at once, but we need to serialize them onto the pins that drive data back across the bus. There are 16 data pins per channel. The 256 bits of our read leave on these pins as PAM4GDDR6X is the GDDR6 standard, with this PAM4 signalling added. symbols: each symbol is one of four voltage levels, carrying two bits, so 256 bits over 16 pins is 16 bits per pin — eight symbols. The clock is sent along a shared wire so that the controller can sample at the right edges.

The way back

These PAM4 bursts are deserialized in the memory controller, and written into the L2 slice’s line. The results run back through the crossbar, back to their SM, and fill their L1 slot. They rendezvous with the record left by their leaving, and their bytes are written into register R4 across all the lanes.

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

The round trip — L1, TLB, crossbar, L2, controller, and back — costs about 255 ns, some 660 cycles. All the while our warp was parked on its barrier. The rest of the chip wasn’t idle though. The sub-partition issued the same loads for another 11 warps, the rest of the SM for another 36, the other SMs for the other 6096. The result is a cacophony of loads, the per-load latency of any one of them lost in 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 in the real kernel: those addresses light up (and miss) in the correct L1 set, then are routed through the crossbar to the correct L2 slice, where they miss, falling through a correctly contended memory controller to a simulated DRAM bank, before returning back through L2, back through L1, and returning their results into the correct register.

Appendix: the probes

Setup

All measurements are on one RTX 4090 (sm_89), with the core clock locked at 2.6 GHz. Cycles come from measured nanoseconds at that frequency. Two main instruments:

A latency chase. To get a latency measurement (especially when that latency changing tells you something about the chip), we run a pointer cycle through a chosen set of lines, hopped 20,000 times, and then measure the mean ns per hop. If the lines we point to fit in a cache level, then they stay resident, and the mean is that level’s hit latency. Because of the steepness of the hierarchy, any loads that overflow to the next level down tend to show up strongly in the average. ld.global.ca (LDG.E…STRONG.SM) for chases at the L1, ld.global.cg (LDG.E…STRONG.GPU) goes past L1. Hit latencies are 15.4 ns at the L1, 127.4 ns at the L2, and 255.4 ns at DRAM.

Hardware counters. To read ncu’s counters reliably you have to take them as slopes over iteration count so fixed overhead cancels. Sector and request counters at the L1 exit port and the L2 side are used to figure out more about the shape of the requests, and a per-slice sector counter helps to give us the L2 slice measurements.

The L1 set function

The 8 bits of the L1 index are the XOR of a fixed subset of the address bits. Written as a bitmask over the address, one basis for those subsets is:

The masks themselves aren’t unique — any invertible combination of these eight describes the same partition.

Page tables and the TLB

The 16-entry TLB is only the first level, but what happens when you miss? A miss refills in about 4.4 ns, and an L2 hit is 127 ns and a VRAM access is 255 ns, so we can’t be going from those. The inference is that it comes from some larger on-chip translation cache.

The cost is flat within 0.1 ns for all the pages the chip can map, and from any SM. More evidence: walking the page tables with nvdebug shows the volatile bit set on every directory entry, so they’re not cached in the normal hierarchy.

The L2 slice function

Measuring which slice owns a line is pretty hard. The L2 is physically indexed, so the probe has to work in device-physical addresses from the page-table walk. Nsight Compute does have a per-slice sector counter, but reports only the min, max, average, and sum across the 36 instances, never the actual slice index.

Even so, the aggregate is enough to tell whether two addresses share a slice. If the two addresses live on the same slice, after loading both, the max counter reports 2, if they’re on different slices the max is 1. You can use this probe to get a representative address that lands on each of the 36 slices.

With the 36 representatives in hand, you can get any new candidate’s slice. If you read the candidate many times alongside all 36, with each of the different addresses read a distinct number of times (say 20001, 20002, … times), the sum of the candidate’s read count and only one of the representatives will match the max counter, and you can figure the slice by inference.

From that, you can produce a table of many physical address-slice pairs. The hard part is going from such a table to a physically plausible function. One tool that helped us a bit was running the same kinds of experiments on two different chips built on the same die: the 4090, and the L40S, which has an extra slice per memory controller.

Here’s one Claude made earlierIt's hard to be sure what's actually in the hardware here, but this is plausible given my limited knowledge. The priors: there's got to be some shared silicon between the L40S and the 4090 (assuming NVIDIA don't ship two completely different functional paths for chips on the same layout but with different amounts of L2 fused off). And the function has to be simple-ish in hardware, i.e. XORs, arithmetic etc. are fair game, but if Claude tries to put in a 4096 entry lookup table you tell it to go try harder.:

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, and it reaches bit 35
    b35 = (1 << 35) if wide else 0

    # stage 1 — which of the 12 controllers: two parities and a mod-3 digit
    P1c = parity(a & 0x76A990400)                # controller parity 1 (narrow; used on both chips)
    P1  = parity(a & (0x76A990400 ^ b35))        # wide form, only needed for the L40S read-out
    P2  = parity(a & 0x2CCF7B000)                # controller parity 2
    A   = ((a >> 15) + 2*parity(a & 0x3C9041000) + parity(a & (0x2882B0800 ^ b35)) + 2) % 3  # mod-3 digit: (a>>15) + 2 corrections

    # stage 2 — which slice inside the controller: a 9-position cyclic counter
    g   = ((a + (1 << 16)) >> 17) % 9            # the counter value, round(a / 2^17) mod 9
    q0  = parity(a & 0x8000)                     # four correction parities
    q1  = parity(a & 0x5985E0500)
    q2  = parity(a & (0x2354E4400 ^ b35))
    q3  = parity(a & 0x3C9041000)
    carry = 1 if q0 + q1 + q2 >= 2 else 0        # q0,q1,q2 as a full adder: the carry (majority)...
    start = (5 + 7*q0 + 5*q1 + 2*q2 + q3 - carry) % 9   # ...sets where the counter starts
    o     = (g - SHIFT[A] - start) % 9           # position within the 9-cycle
    Lf    = 2 if (q0 ^ q1 ^ q2) == 0 else 1      # ...and their XOR sets where it splits
    return P1c, P1, P2, A, q2, o // 3, (1 if (o % 3) >= Lf else 0)   # d = o // 3, u = the split bit

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 live, read (d, u) as three arcs of Z/9
        base = 2 if d == 0 else (1 if (d == 1 and u == 0) else 0)
        B = ((1 - base) % 3 if q2 else base) % 3 # q2 flips the arc order
        B = (B + OFFSET[A]) % 3                   # per-controller offset
        return controller * 3 + B
    if N == 48:                                  # L40S: 4 slices live, read u as two index bits
        i0, i1 = P1 ^ q2 ^ u, P1 ^ P2 ^ u
        return controller * 4 + 2*i0 + i1
    raise ValueError("N must be 36 or 48")

Whilst it is very hard to find such a function, it’s very easy to tell if you’ve found one that works. Drawing 8,192 L2-resident lines from exactly k predicted slices:

The L2 set index and geometry

Once the slice function pins addresses to a single slice, you can do the same eviction-set archaeology on that slice, to figure out the structure, which tells you that it’s 16 way set-associative (a chase with 17 elements thrashes, but one with 16 doesn’t).

The set index within a slice is the same kind of parity function as the L1’s — ten bits, with the same (a >> 15) mod 9 nonlinearity in the top bit. Unfortunately, the masks involved differ depending on the slice. For one slice:

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

def set_index(a):                                # within one slice
    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]))

It has some properties that let you sense-check it. For example: a contiguous 72MiB fills each slot in each slice without thrashing anything, as you’d expect.

DRAM refresh

DRAM cells leak charge and so have to be periodically refreshed, which makes some kinds of timing probes harder. You can see it by running a dependent chase that writes each hop’s timing into shared memory. Most DRAM accesses come back at the usual latency, but a small share take longer, spread evenly out to a hard ceiling about 210 ns higher than usual. An evenly spaced run like that is the signature of a fixed length stall. The stall is ~210 ns. About 2% of accesses hit one. It doesn’t hit the whole chip at once — it’s more local than that — but I couldn’t tell what the unit was.