Anatomy of a GPU Memory Read

Originally published by Fergus Finn, Founder Member of Technical Staff at Doubleword.

In a previous exploration, we tracked a vector-addition kernel—defined by the operation $c[i] = a[i] + b[i]$—from the nvcc compiler down to the execution of warps. While we covered the launch process, several internal hardware mechanisms remained unexplored.

This article fills those gaps. We will trace the lifecycle of a critical SASS instruction (a global load) as it traverses the hardware of an RTX 4090. This level of reverse engineering is driven by a need for performance optimization; as noted in the Citadel microbenchmarking paper, these granular details are where the real gains are found.

NVIDIA provides extensive documentation $\rightarrow$ Actually, NVIDIA provides very little detail at this level, so we rely on hardware-based timing experiments to map the path.

The Target: A Simple Vector Add

The CUDA kernel under the microscope is straightforward:

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

When compiled to SASS, the logic for loading elements of vector b looks like this:

/*0080*/ IMAD.WIDE R4 , R6 , R7 , c [0x0][0x168] ; // Calculate address for b[i]
/*00a0*/ LDG.E R4 , [R4.64] ;                      // Load b[i] from global memory

The LDG.E instruction moves data from global memory into a register. To fulfill this single request, the hardware must coordinate:

• Four 32-byte sectors.
• One full cache line.
• One address translation.
• A crossbar traversal.
• One of 36 L2 slices.
• (On a miss) An activation and four column reads from a DRAM chip.

---

The Execution Context

Our warp resides on one of the four sub-partitions of a Streaming Multiprocessor (SM), sharing space with eleven other resident warps. The sub-partition scheduler selects one eligible warp per cycle to issue an instruction across all 32 lanes.

In our scenario, the warp is selected twice:

1. First, for the IMAD.WIDE (to calculate the address).
2. A few cycles later, for the LDG (once the addresses are ready in registers R4 and R5).

The Hardware Path

---

Phase 1: From Warp to L1 Cache

The instruction LDG.E R4, [R4.64] requests a 32-bit value from a 64-bit address. Because registers are only 32 bits wide, the .64 notation indicates that the address is split across R4 and R5.

Register Retrieval

To begin, the hardware must fetch the addresses from the register file:

• Parallelism: One row of the register file provides R4 for all 32 lanes simultaneously.
• Operand Collector: Reads are staged here to prevent bottlenecks. Since a register bank can only serve one read per cycle, the collector manages conflicts.
• Banking: There are two banks, determined by the register number's lowest bit. Consequently, any two adjacent registers (like R4 and R5) are guaranteed to be in different banks.

This process yields 256 bytes (32 lanes $\times$ 8 bytes per address).

Timing Insight: Register Access Cost The address read adds $\approx 1$ cycle. For comparison, a shared-memory load using a register-based address takes 24 cycles from issue to first use, whereas one using an immediate address takes 23 cycles.

The Load/Store Unit (LSU) and Coalescing

Once addresses are resolved, the instruction moves to the LSU. The LSU handles:

1. Address Arithmetic: Adding immediate offsets (though none are present here).
2. Scope Resolution: LDG explicitly targets the global memory window.

The LSU then forwards the opcode, a 32-bit active lane mask, the computed addresses, and the destination register.

The Coalescing Challenge: While each lane asks for 4 bytes, the L1 cache is organized into 32-byte sectors. The warp coalescer determines the most efficient way to group these requests. In this case, it identifies that the 32 lanes are requesting 128 contiguous bytes, resulting in 4 contiguous sector requests.

---

Phase 2: Entering the L1 Cache

The request for four 32-byte sectors is sent to the L1 cache, which is further organized into 128-byte lines. Since our request covers exactly 128 bytes, it maps to a single cache line.

L1 Architecture

The L1 cache on the RTX 4090 is 4-way set-associative. This means:

• Sets: The cache is divided into sets.
• Slots: Each set contains four slots.
• Tags: Each slot has a tag to identify the stored line.

The lookup process compares the requested line's tag against the four tags in the designated set.

The Result

If a tag matches and the required sectors are present, the data is returned immediately. However, since this is our first access, the request misses and must proceed deeper into the memory hierarchy.

Performance Metric:

• L1 Hit Latency: $\approx 15.4\text{ ns}$ (roughly 40 cycles).