GoKawiilOur 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 4090 We 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 b The 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 .
warp coalescer L1 15 ns TLB crossbar L2 127 ns controller DRAM 255 ns SM die board
... continue reading