What happens when you run a CUDA kernel

Here’s a simple CUDA program. It adds two vectors. __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]; } int main() { int n = 1 << 20; // a million floats (1,048,576) size_t bytes = n * sizeof(float); 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.0f; float *da, *db, *dc; cudaMalloc(&da, bytes); cudaMalloc(&db, bytes); cudaMalloc(&dc, bytes); cudaMemcpy(da, a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(db, b, bytes, cudaMemcpyHostToDevice); vadd<<<4096, 256>>>(da, db, dc, n); // 4096 * 256 = n threads, one per float cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost); printf("c[0]=%f c[n-1]=%f\n", c[0], c[n-1]); } Compiled for an RTX 4090, and launched, it does correctly work out that 1+1=21+1=2, a million timesI didn’t check all of them.. $ nvcc -arch=sm_89 -o vadd vadd.cu && ./vadd c[0]=2.000000 c[n-1]=2.000000 Telling you that involved tens of millions of CPU instructions, a couple of device files, nine hundred ioctls, and one memory-mapped doorbell register. In this post, we’ll follow this one kernel from the code down to the warps, and back up to the answerAn aside, this post is an instance of the ‘legibility transition’ that agents have engendered.

There really is very little about computers you can’t find out with curiosity and (machine-enhanced) persistence. An interesting discussion of the implications of legibility for what AI can help us to know here.. Compiling our program with nvcc§ We ought to start with how to turn this CUDA program into something that the device can actually read. To do that we need a compiler. Really, we need many compilers. nvcc is a driver program that runs several other compilers and combines their output. If you pass --keep it leaves the whole pipeline on disk for you to read: $ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ls ... vadd.ptx # device code as PTX (from cicc) vadd.sm_89.cubin # device code as SASS (from ptxas) vadd.fatbin # cubin + PTX, bundled (from fatbinary) vadd.cudafe1.stub.c # host launch stub + kernel registration vadd.o # final host object, fatbin embedded ... The host code goes to your host compiler. The device code (vadd) takes more steps: cicc, an LLVM-based compiler, turns it into PTX, and then ptxas turns the PTX into SASS. PTX is a virtual ISA. It has infinitely many typed registers, and no notion of how many of them the hardware actually has. Here is the (elided) body of vadd in PTX: $ cat vadd.ptx ... mad.lo.s32 %r1, %r3, %r4, %r5; // set register r1 to ctaid*ntid + tid setp.ge.s32 %p1, %r1, %r2; // set predicate p1 if i >= n @%p1 bra $L__BB0_2; // if out of bounds, skip to exit cvta.to.global.u64 %rd4, %rd1; // convert generic pointer %rd1 to a global address, store in %rd4 mul.wide.s32 %rd5, %r1, 4; // multiply r1 by 4, store the result in %rd5 add.s64 %rd6,

%rd4, %rd5; // add %rd4, %rd5, result in %rd6 ld.global.f32 %f2, [%rd6]; // load a[i] into %f2 ... add.f32 %f3, %f2, %f1; // add %f1 and %f2, result in %f3 st.global.f32 [%rd10], %f3; // store c[i] = ... in global memory The virtual registers look like %rd1–%rd10, %f1–%f3The prefix is the type: %r is a 32-bit integer, %rd a 64-bit one, %f a 32-bit float, %p a one-bit predicate.. PTX is more ‘longhand’ than you might expect. For example, forming one address in %rd6 takes three PTX instructions. This happens because PTX is device agnostic. Why three? CUDA pointers are “generic” by default, meaning they could name global, shared, or local memory. cvta.to.global asserts the pointer lives in the global window, so a cheaper ld.global can be used later. mul.wide.s32 then turns the index i into a byte offset by multiplying by 4 (sizeof(float)) and widening 32→64 bits in one step. add.s64 adds that to the base pointer. Next, ptxas transforms our PTX, which is device agnostic, into the SASS for your architecture, which isn’t. The SASS it emits looks different: $ cuobjdump -sass vadd /*0000*/ MOV R1, c[0x0][0x28] ; // set up the stack pointer (ABI; unused here) /*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 ; // if so, exit /*0060*/ MOV R7, 0x4 ; // load literal 4 (sizeof(float)) into R7 as multiplier /*0070*/ ULDC.64 UR4, c[0x0][0x118] ; // uniform load of a driver-provided system value /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*0090*/ IMAD.WIDE R2, R6, R7, c[0x0][0x160] ; // &a[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i] /*00b0*/ LDG.E R3, [R2.64] ; // a[i] /*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] ; // &c[i] /*00d0*/ FADD R9, R4, R3 ; // a[i] + b[i] /*00e0*/ STG.E [R6.64], R9 ; // c[i] = ... /*00f0*/ EXIT ; What the S2R lines are doing S2R is “special register to register”: it copies a special register the hardware maintains per thread — here SR_CTAID.X (the block’s index, blockIdx.x) and SR_TID.X (the lane’s index within the block, threadIdx.x) — into an ordinary register so IMAD can do arithmetic on it. Ten-odd virtual registers have collapsed onto seven real onesncu reports launch__registers_per_thread = 16. The disassembly only names up to R9, but the allocator reserves a few more for the ABI and alignment.. The two mul.wide plus add sequences have fused into a single IMAD.WIDE.

Finally, that fatbin is nested in the host executable, where readelf -S finds it occupying its own sections: $ readelf -S vadd ... [18] .nv_fatbin PROGBITS ... [19] __nv_module_id PROGBITS ... [29] .nvFatBinSegment PROGBITS ... ... The vadd binary that nvcc spits out is a single executable containing host code, a complete ELF object containing the Ada SASS, and a copy of the PTX. Because PTX is verbose plain text, nvcc compresses it by default to keep the binary size small; the driver will only decompress and JIT-compile it if the binary is run on an architecture that the pre-compiled SASS doesn’t cover. How the host triggers the GPU§ The compiled GPU machine code is now sitting inert inside the .nv_fatbin section of our ./vadd executable. When you launch the program on the host, we have to bridge two worlds: the host CPU, and the GPU sitting across the PCIe bus. To set up a host binary that knows how to cross the bridge, the frontend compiler (cudafe++) inserts a hidden constructor into your code, running before the main function starts. Its job is to register our embedded fatbinary with the CUDA runtime and record a mapping that the runtime will later use: associating the host-side function pointer vadd with the compiled device kernel’s mangled name in the fatbin. When the compiler encounters vadd<<<4096, 256>>>(da, db, dc, n), it replaces that high-level expression with a generated host launch stub. This stub packs our kernel arguments into a buffer in host memory. The pointers da, db, dc and the integer n are aligned at byte offsets 0, 8, 16, and 24 These offsets are the constant bank offsets 0x160, 0x168, 0x170, and 0x178 that we saw our SASS machine code reading from constant bank 0 earlier.: //

from vadd.cudafe1.stub.c void __device_stub__Z4vaddPKfS0_Pfi(const float *__par0, const float *__par1, float *__par2, int __par3) { __cudaLaunchPrologue(4); __cudaSetupArgSimple(__par0, 0UL); // arg buffer offset 0 __cudaSetupArgSimple(__par1, 8UL); // offset 8 __cudaSetupArgSimple(__par2, 16UL); // offset 16 __cudaSetupArgSimple(__par3, 24UL); // offset 24 __cudaLaunch((char*)(void(*)(const float*, const float*, float*, int))vadd); } Once the arguments are packed, the stub calls __cudaLaunch, passing it the memory address of the host-side dummy vadd function. Because this host function is just an empty shell on the CPU, its host memory address serves as a lookup key. The runtime queries its registration table with this address to find the corresponding device-side symbol name, and then crosses the boundary into the closed-source user-mode driver (libcuda.so.1)The usermode bit of the driver comes with the GPU’s kernel driver, not with the CUDA toolkit: the libcuda.so.1 from the strace resolves to libcuda.so.590.48.01, the driver release on this machine. to initiate the launch of that kernel. The runtime opens this driver dynamically on the first GPU call in our program, which we can catch using strace: $ strace -f -e trace=openat ./vadd ... openat(..., "/lib/x86_64-linux-gnu/libcuda.so.1", O_RDONLY|O_CLOEXEC) = 3 ... When this first call is performed, a ‘context’ is created, containing all the infrastructure the driver needs to talk to the device, including the channel through which the CPU speaks to the GPU. We’ll talk more about that in the next section. At this stage, the compiled machine code still hasn’t reached the GPU. Since CUDA 12.2, module loading is lazy by defaultControlled by CUDA_MODULE_LOADING.