GPU Architecture · Execution & Operations

The launch path: from <<<>>> to a warp on an SM

One kernel launch, traced end to end — what your compiler baked in, what the driver does, and how the work actually crosses into the GPU.

You've built two maps: the physical hardware (SMs, warps, schedulers, the memory tiers) and the logical CUDA model (kernel, grid, block, thread). This lesson adds a third thing that sits between them — the plumbing that carries a launch from your host process, through the driver, into the hardware you already understand. When a kernel misbehaves, the bug is often in this layer, not in your math. So let's follow a single launch the whole way down.

It starts at compile time, not run time

nvcc is not really a compiler — it's a compiler driver that runs your .cu file down two paths and staples the results together:¹

Decoding the toolchain vocabulary

Those two paths sling a lot of acronyms. Pin them down once — this is the vocabulary the whole ecosystem, and every compiler error, is written in:

nvcc

The driver you invoke. It orchestrates the two paths and calls the tools below; it doesn't compile device code itself.

PTX

Parallel Thread Execution — a virtual instruction set: portable, text, human-readable assembly. The GPU never runs it directly; it's an intermediate stage, roughly "the GPU's LLVM IR." Targets a virtual architecture, compute_XX.

ptxas

The PTX assembler — the optimizing back-end that turns PTX into real machine code for one specific GPU. Register allocation and most device-code optimization happen here, not in the front end.

SASS

The real, binary machine instructions an SM executes (the name is historical — "Streaming ASSembler"). This is the "device machine code" referred to throughout. Targets a real architecture, sm_XX.

compute_XX
sm_XX

The target architecture. compute_XX = virtual (what PTX targets); sm_XX = real (what SASS targets). The number XX is the GPU's compute capability — its hardware feature/ISA generation.

cubin

A CUDA binary: the ELF file ptxas emits, holding the SASS plus metadata for one architecture.

cuobjdump
nvdisasm

The inspector tools — they crack open a binary or cubin to show you the embedded PTX and disassembled SASS. (You'll use cuobjdump in the exercise below.)

The consequence is worth sitting with: your CPU executable literally contains GPU machine code inside it. On Linux/ELF it lands in a section named .nv_fatbin. That's the first piece of "metadata that travels with the kernel" — and there's more riding along inside each cubin: the SASS, plus a .nv.info section recording each kernel's resource needs (how many registers, how much shared memory, the parameter layout).²

One more thing nvcc plants in the host binary: a hidden registration constructor that runs before main(). It hands the fatbinary to the CUDA runtime and, for every kernel, records a mapping from a host-side stub function to the kernel's mangled device name. That table is how a launch later finds the right GPU code. (Documented that this happens; the exact symbol names — __cudaRegisterFatBinary, __cudaRegisterFunction — are visible in the generated stubs but are internal, not a public contract.)²

# See it for yourself — the GPU code is right there in your binary:
cuobjdump -sass ./mat2d          # the SASS for each embedded kernel
cuobjdump -lelf ./mat2d          # list the cubin images inside the fatbin
readelf -S ./mat2d | grep nv_fatbin   # the section holding the device code
nm ./mat2d | grep -i cudaRegister     # the (internal) registration symbols

First touch wakes the device

Nothing GPU-side has happened yet. The first CUDA call in your program that needs the device triggers lazy initialization: the runtime creates the device's primary context — think of a context as a process on the GPU: its own address space, its allocations, its loaded modules, its streams. This is also when any PTX gets JIT-compiled and modules are loaded into device memory.³

What <<<>>> actually compiles to

The triple-chevron is pure syntactic sugar. For myKernel<<<grid, block, shmem, stream>>>(a, b), nvcc emits roughly:⁴

// 1. push the launch config (grid, block, shared-mem bytes, stream) onto a TLS slot
__cudaPushCallConfiguration(grid, block, shmem, stream);
// 2. call the generated host STUB for this kernel, which:
//    - packs &a, &b into an array of argument pointers
//    - pops the config back
//    - makes the one real runtime call:
cudaLaunchKernel(func, grid, block, args, shmem, stream);

Two things here repay attention:

• func is a host address used only as a lookup key — it indexes the registration table from earlier to find the actual device code. It is not a pointer to GPU memory.
• Your arguments are copied, by value, into a small parameter buffer that ships with the launch and is delivered to the kernel through constant memory (the PTX .param space you met in Lesson 5). There's a hard size cap on this buffer: 32,764 bytes on your GPU (Volta and newer, CUDA 12.1+); it was just 4 KB before. Pass a giant struct by value and you'll hit it.⁴

And critically: cudaLaunchKernel is asynchronous. It hands the work off and returns to your CPU thread immediately — usually before the kernel has even started running, let alone finished. (That single fact drives all of Lesson 8: completion, timing, and error reporting.)⁵

Crossing into the GPU

Now the work leaves your process. The user-mode driver writes the launch as GPU commands into a pushbuffer, links that into a per-context ring buffer (the GPFIFO), and rings a doorbell — a single write to a register mapped into your process's address space. No system call per launch; that's what keeps launch overhead down to microseconds. On the GPU, the host interface sees the doorbell, reads the ring, and the GigaThread Engine — the global work distributor — begins handing your blocks to SMs, exactly as Lesson 3 described. Here is the whole path on one page:

1. host process
   Your executable — already carries the kernel's SASS in an embedded fatbinary, and registered every kernel before main().
2. CUDA runtime
   First CUDA call creates the context, JIT-compiles PTX if needed, loads modules into device memory.
3. runtime + UMD
   <<<>>> → stub → cudaLaunchKernel. Args copied into a parameter buffer. Returns immediately.
4. user-mode driver
   Writes GPU commands to a pushbuffer, appends a GPFIFO ring entry, rings the doorbell — no syscall.
PCIe
5. host interface
   Sees the doorbell, DMA-fetches the commands from the ring.
6. GigaThread engine
   Global work distributor assigns each block to an SM (indivisible, whole block to one SM).
7. SM warp schedulers
   Issue the block's warps of 32 threads — the launch is now running (Lessons 2–3).

Notice the arc closes exactly where the compute lessons began. Everything you learned about occupancy, warps, and scheduling picks up at the last two rows; this lesson is the four rows above them that you'd never seen.

write GPU commands   → pushbuffer    // ordinary RAM store
advance ring PUT ptr → GPFIFO        // ordinary RAM store
write channel id     → doorbell reg  // MMIO store → pokes the GPU

The honesty seam: what's documented, what's reverse-engineered

This matters for troubleshooting — it tells you which behaviors you can rely on and which are implementation details that could shift. NVIDIA documents the model down to cudaLaunchKernel and the PTX parameter space precisely. Below that, the pushbuffer/doorbell mechanics are pieced together from NVIDIA's open-source kernel modules, published hardware manuals, and community work — real, but not an API contract:

Check yourself

From memory — one click locks each answer.

Notes & citations

1. CUDA Compiler Driver NVCC, "The CUDA Compilation Trajectory," and CUDA C++ Programming Guide §3.1 (PTX vs cubin, fatbinary, just-in-time compilation of PTX by the driver). docs.nvidia.com/cuda/cuda-compiler-driver-nvcc
2. CUDA Binary Utilities (cubin is ELF; .text, .nv.info, symbol sections). That device code is embedded and inspected by the runtime is documented in the NVCC guide; the exact .nv_fatbin section name and the __cudaRegister* stub symbols are community-verified via nvcc --keep / readelf, not a public API. docs.nvidia.com/cuda/cuda-binary-utilities
3. CUDA C++ Programming Guide, "Runtime Initialization" — the primary context is created lazily on the first runtime call needing it (eagerly at cudaSetDevice since CUDA 12.0); a context is "analogous to a CPU process." docs.nvidia.com/cuda/cuda-c-programming-guide
4. CUDA Runtime API, "Execution Control" — cudaLaunchKernel(func, gridDim, blockDim, args, sharedMem, stream), where args is an array of pointers to the actual parameters. Large-parameter limit (32,764 bytes on sm_70+, CUDA 12.1 / R530): CUDA 12.1 Supports Large Kernel Parameters. The __cudaPush/PopCallConfiguration desugaring is visible in generated stubs (nvcc --keep), not separately documented.
5. CUDA C++ Programming Guide, "Asynchronous Concurrent Execution" — kernel launches are asynchronous and return control to the host immediately. docs.nvidia.com
6. NVIDIA CUDA on WSL User Guide and Microsoft "DirectX ❤ Linux" — the WSL2 path: stubbed libcuda.so → /dev/dxg → Linux dxgkrnl → VM bus → Windows host driver (GPU paravirtualization). The pushbuffer/GPFIFO/doorbell model is drawn from NVIDIA's open-gpu-doc hardware manuals and open-gpu-kernel-modules. docs.nvidia.com/cuda/wsl-user-guide