Our previous post followed a vector-add kernel from nvcc down to the warps. This time, we follow a critical SASS instruction (a global load) through the hardware of an RTX4090. The CUDA kernel investigated has two lines: __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];}. The compiled SASS shows: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]. These instructions load elements of vector b from global memory into registers.

One LDG.E asks for four bytes in each of 32 lanes. Serving it requires four 32-byte sectors, one cache line, one address translation, a crossbar crossing, one of thirty-six L2 slices, and when missing everywhere, an activate and four column reads at a DRAM chip. Our warp lives on one of the SM's four sub-partitions alongside eleven other resident warps. The instruction issues to the load/store unit (LSU), which sends the opcode, active lane mask, computed addresses, and register number to the coalescer.

The coalescer determines that 4 contiguous sector requests are needed for the 128 bytes the warp requested. Entering the L1 cache, the request is sent to the virtually addressed, 4-way set-associative cache. The lookup compares tags in sets of four slots. Since this is a first-time load, it misses and descends further. An L1 hit returns in about 15.4 ns — 40 cycles. Virtual memory requires translation before leaving the SM. The SM's TLB holds sixteen recent translations. Misses cost about 4.4 ns — eleven cycles.

After translation, one request per 128-byte line proceeds across the crossbar to one of 36 2 MiB L2 slices, picked by a function of its physical address. All slices can serve in parallel, giving aggregate bandwidth 36x that of a single slice.