What happens when a GPU reads memory | Doubleword

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. What the register retrieve costs The address read adds at most one cycle. A shared-memory load taking its address from a register takes 24 cycles from issue to first use, and the same load with the address as an immediate takes 23. (LDG can’t take an immediate). 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 for1. 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 slots2, 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 program3. 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 engineer4. 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 done5. Because we’re loading all of our data for the first time, our request misses, and must descend further into the memory system. How much does an L1 hit cost An L1 hit returns in about 15.4 ns — 40 cycles. The number comes from one thread chasing a dependent chain through a random permutation of L1-resident lines, with the latency chase. 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 SM6. 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 assignment7. 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 warps8. The very first load will miss in this TLB. What translation costs We can’t see any cost to hitting the TLB in any of the probes we have. Misses cost about 4.4 ns — eleven cycles. The same refill cost holds within 0.1 ns across all the pages this chip can map, and from any SM, so the next level of the translation cache is universal, and very cheap. 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 marked9. 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 address10. 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-associative11. 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 chip12. Our request gets handed over to that controller. What does this cost An L2 hit costs about 127 ns — some 330 cycles. Each SM can hand the crossbar up to two line-requests per cycle, and the 36 slices serve independently. The exit-port counter is l1tex__m_l1tex2xbar_req_cycles_active. 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). GDDR6Xchannel 0 · 1 GiBchannel 1 · 1 GiBone bank · 65,536 rows of 1 KiB⋯one row · 1 KiB · 32 columns of 32 B 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 row13. So, to serve our load, the memory controller must first send one activate, and then four reads14. 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