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.

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