Anatomy of a CUDA Kernel Execution

What actually happens under the hood when you trigger a CUDA kernel? To explore this, let's look at a basic vector addition example.

The Source Code: vadd.cu

The following program performs a simple element-wise addition of two arrays.

__global__ void vadd ( const float* a, const float* b, float* c, int n) { 
    // Calculate the unique global index for this thread
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    if (i < n) 
        c[i] = a[i] + b[i]; 
} 

int main () { 
    int n = 1 << 20 ; // 1,048,576 elements
    size_t bytes = n * sizeof ( float ); 
    
    // Host allocation
    float * a = ( float* ) malloc (bytes), * b = ( float* ) malloc (bytes), * c = ( float* ) malloc (bytes); 
    for ( int i = 0 ; i < n; i ++ ) a[i] = b[i] = 1.0 f ; 
    
    // Device allocation
    float * da, * db, * dc; 
    cudaMalloc ( da, bytes); 
    cudaMalloc ( db, bytes); 
    cudaMalloc ( dc, bytes); 
    
    // Data transfer: Host -> Device
    cudaMemcpy (da, a, bytes, cudaMemcpyHostToDevice); 
    cudaMemcpy (db, b, bytes, cudaMemcpyHostToDevice); 
    
    // Kernel Launch: 4096 blocks * 256 threads = 1,048,576 threads
    vadd<<<4096, 256>>>(da, db, dc, n); 
    
    // Data transfer: Device -> Host
    cudaMemcpy (c, dc, bytes, cudaMemcpyDeviceToHost); 
    printf ( "c[0]= %f c[n-1]= %f\n" , c[ 0 ], c[n - 1 ]); 
}

When compiled for an RTX 4090, the result is a mathematically trivial but computationally massive operation: $\text{1 + 1 = 2}$, repeated $1,048,576$ times.

$ nvcc -arch=sm_89 -o vadd vadd.cu
./vadd
c[0]=2.000000 c[n-1]=2.000000

While the output is simple, the process is not. This execution involves tens of millions of CPU instructions, several device files, roughly 900 ioctls, and a specific memory-mapped doorbell register.

A Note on Legibility: This analysis is a product of the "legibility transition" enabled by AI agents. With enough curiosity and machine-assisted persistence, almost any aspect of modern computing can be uncovered.

---

🛠️ The Compilation Pipeline

To understand how the device reads this code, we must look at nvcc. Rather than being a single compiler, nvcc is a driver that orchestrates several tools. By using the --keep flag, we can inspect the intermediate artifacts:

$ nvcc --keep -arch=sm_89 -o vadd vadd.cu
ls ...
# vadd.ptx             <- Device code as PTX (via cicc)
# vadd.sm_89.cubin     <- Device code as SASS (via ptxas)
# vadd.fatbin          <- Bundled cubin + PTX (via fatbinary)
# vadd.cudafe1.stub.c  <- Host launch stub & registration
# vadd.o               <- Final host object with embedded fatbin

The Transformation Flow

1. PTX (Parallel Thread Execution)

The cicc compiler produces PTX, a low-level virtual ISA. PTX assumes an infinite supply of typed registers and ignores the physical constraints of the hardware.

Example PTX snippet:

mad.lo.s32 %r1, %r3, %r4, %r5 ; // i = blockIdx.x * blockDim.x + threadIdx.x
setp .ge.s32 %p1, %r1, %r2 ;    // predicate p1 = (i >= n)
@%p1 bra $L__BB0_2 ;            // if p1 is true, jump to exit
cvta.to.global.u64 %rd4, %rd1 ; // convert generic pointer to global address
mul .wide.s32 %rd5, %r1, 4 ;    // i * sizeof(float)
add .s64 %rd6, %rd4, %rd5 ;     // final address for a[i]
ld.global.f32 %f2, [%rd6] ;     // load a[i]
...
add .f32 %f3, %f2, %f1 ;         // a[i] + b[i]
st .global.f32 [%rd10], %f3 ;    // store in c[i]

PTX Register Types:

Key Nuances in PTX:

• Generic Pointers: CUDA pointers are generic by default. The cvta.to.global instruction explicitly asserts the pointer is in global memory, allowing the use of the more efficient ld.global.
• Address Calculation: Converting an index $i$ to a byte offset requires mul.wide.s32 to multiply by 4 (the size of a float) and widen the result from 32 to 64 bits.

2. SASS (Streaming Assembler)

Next, ptxas converts the device-agnostic PTX into SASS, which is specific to the GPU architecture (e.g., sm_89 for the RTX 4090).

$ cuobjdump -sass vadd
/*0000*/ MOV R1 , c [0x0][0x28] ; // Setup stack pointer
/*0010*/ S2R R6 , SR_CTAID.X ;    // R6 = blockIdx.x
/*0020*/ S2R R3 , SR_TID.X ;      // R3 = threadIdx.x
/*0030*/ IMAD R6 , R6 , c [0x0][0x0] , R3 ; // i = ctaid*ntid + tid
/*0040*/ ISETP.GE.AND P0 , PT , R6 , c [0x0][0x178] , PT ; // P0 = (i >= n)
/*0050*/ @P0 EXIT ;               // Exit if out of bounds
/*0060*/ MOV R7 , 0x4 ;           // sizeof(float)
/*0070*/ ULDC.64 UR4 , c [0x0][0x118] ; // Load system value
/*0080*/ IMAD.WIDE R4 , R6 , R7 , c [0x0][0x168] ; // Address of b[i]
/*0090*/ IMAD.WIDE R2 , R6 , R7 , c [0x0][0x160] ; // Address of a[i]
/*00a0*/ LDG.E R4 , [R4.64] ;     // Load b[i]
/*00b0*/ LDG.E R3 , [R2.64] ;     // Load a[i]
/*00c0*/ IMAD.WIDE R6 , R6 , R7 , c [0x0][0x170] ; // Address of c[i]
/*00d0*/ FADD R9 , R4 , R3 ;      // a[i] + b[i]
/*00e0*/ STG.E [R6.64] , R9 ;     // Store result
/*00f0*/ EXIT ;

Observations on SASS:

• S2R Instructions: "Special Register to Register" moves hardware-maintained values (like SR_CTAID.X for blockIdx.x) into general-purpose registers for arithmetic.
• Register Compression: Ten-odd virtual registers $\rightarrow$ Seven real registers. While ncu reports 16 registers per thread, the disassembly only uses up to R9 (the rest are for ABI/alignment).
• Instruction Fusion: The separate multiply and add steps from PTX are fused into a single IMAD.WIDE instruction.

---

📦 Kernel Arguments and Memory

The operands appearing as c [0x0][...] refer to Constant Bank 0. This is a small, driver-managed memory region.

How Arguments are Handled

When a kernel is launched, the driver provides a structure called the QMD (Queue Management Descriptor) to the GPU. This fills Constant Bank 0 with:

1. The pointers a, b, and c.
2. The size n.
3. The launch geometry (grid and block dimensions).

Why Constant Memory? Because these arguments are broadcast reads. Every single thread in the entire grid requires the exact same pointers and size; constant memory is optimized for this specific access pattern.

Summary Checklist of Execution

• nvcc compiles source to PTX.
• ptxas compiles PTX to SASS.
• Host allocates memory and copies data to device.
• Driver populates Constant Bank 0 via QMD.
• GPU schedules warps and executes SASS.
• Results are copied back to host.