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) ⚡

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.

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.

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.

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.

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¹⁴.

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.

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.

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")

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]))