← Chapter 21: Why RAM isn't visible to the GPU Chapter 23: What syncthreads protects →
date 16 / jul / 26
mtwtfss
THE HARDWARE HIERARCHY — Chapter 22 index home

Kernel launch overhead - the hidden tax

*you don't send a cargo ship to deliver one letter*

Launching work on a GPU is not free, and this surprises people because they think of it like calling a function. A normal function call in your code costs less than a nanosecond, a quick jump and you are running. Launching a piece of work on the GPU costs somewhere around 3 to 10 microseconds before the first instruction even executes. That is roughly five thousand times heavier. So if you fire up the GPU just to add two numbers, it's like dispatching a cargo ship to carry one letter across town. The ship is mighty, but the paperwork to launch it dwarfs the errand.

LAUNCHING IS NOT A FUNCTION CALLVSCPU FUNCTION CALLfoo(a, b);JUMPRun< 1 nsA quick jump insidethe same processor.GPU KERNEL LAUNCHkernel<<<grid, block>>>()Many steps before thefirst instruction runs!~ 3 - 10 µs(5000× heavier)Like sending a cargo shipto carry one letter.A GPU launch carries a fixed entry fee of ~3-10 microseconds.Fire it for something tiny, and the fee is all you pay for.

Why The Entry Fee Exists

Start with why that entry fee even exists. When you call a function on the CPU, the processor just jumps to another spot in its own memory and keeps going. Instant. But launching GPU work is not a jump, it's a formal handoff between two separate processors that don't share memory, more like a request sent across a network than a local call. And that handoff has stages that cannot be skipped.

First, the driver wakes up and validates the request, checking that the work dimensions you asked for are actually legal for this particular GPU. Second, it bundles all your arguments and configuration into a command package in CPU memory. Third, it pushes that package across the bridge to the GPU and rings a doorbell, a signal that work is waiting. Fourth, the GPU's own command processor wakes, reads the package, reserves the registers and memory the work needs, and maps it onto its processing units. Only after all of that does the first thread actually run. Every one of those steps costs real time, and none of it is the math you care about. It's all just the protocol of one processor formally asking another to do something.

WHY THE ENTRY FEE EXISTS: THE HANDOFFLike one processor mailing a request across a network to another.CPU (Host)CPUSystem RAMGPU (Device)GPUVRAM1Driver validates the request2Build the command package3Send over PCIe, ring doorbell4GPU command processor wakes5Reserve registers & memory6Map the work onto the SMs7First thread finally executesNot a jump - it is a formal cross-processor handshake.Every step costs real time, and none of it is your math.

↳ So the question that bothers people. If the GPU is sitting there idle and ready, why can't the CPU just poke it and say "go"? Why this whole ceremony? Because the two are genuinely separate machines with separate memory, and starting work across that boundary requires the protocol. NVIDIA has shaved this down over the years and built tools to amortize it, but the basic physics of one chip writing commands that a different chip must parse and schedule cannot be reduced to zero. Those few microseconds are the practical floor.

THE BRIDGE IS THE BOTTLENECKInside GPU (on-die)Superhighway> 1000 GB/sThe bridge is15 - 30×NARROWERPCIe bridge (off-chip)Narrow Bridge~ 20 - 60 GB/sEvery byte that crosses the bridge is expensive.That narrow bridge, not the GPU, sets the real pace.

The Consequence: Losing to a Single CPU Core

Now the consequence that catches everyone, the GPU losing to a single CPU core. It feels impossible. The GPU has thousands of cores, the CPU has a handful, so surely on any task the GPU wins. But on a tiny task it loses badly, for three reasons.

First, clock speed. A CPU core ticks much faster than a GPU core. For a job that's just a few steps done in sequence, the faster ticking finishes sooner, plain and simple. Second, the CPU is built for exactly this. It's loaded with big caches and clever machinery that fast-tracks small sequential work, while a GPU core is lean and not designed to rush a single thread. Third, and decisively, the CPU pays no entry fee. It starts a tiny task in a fraction of a nanosecond. The GPU pays its multi-microsecond launch tax before doing anything.

WHY YOU CAN'T JUST USE THE FAST MEMORYLow Latency (CPU RAM)A Vending Machine(1 snack at a time)Great for one person.Hopeless for 10,000.High Throughput (VRAM)A Dump Truck(10,000 snacks at once)VRAMSlower to arrive,but it feeds everyone.KEY TAKEAWAYYou can't buildone memory thatis BOTH lowlatency AND highthroughput.Physics says no.

So for any job small enough that the actual work is smaller than the launch tax, the CPU wins by default. The rough rule people use is that GPU work should run for at least about a millisecond to be worth launching at all. Anything in the microsecond range and the tax eats you alive.

WHY THE GPU LOSES TO ONE CPU CORE ON SMALL WORK1. CLOCK SPEEDCPU core clock>> GPU core clockFew steps?CPU finishes sooner.CPUGPU2. ARCHITECTURECPU: big caches,branch predictors,built for small,sequential work.GPU: many cores,lean per-core,not built to rusha single thread.CPU COREGPU CORE3. ENTRY FEECPU: ~0 nsGPU: ~3-10 µsIf work < fee,CPU wins by default.CPUGPURULE OF THUMB: GPU work should run for at least~1 millisecond before it is worth launching at all.

Scale: The Only Real Defense

The cleanest way to feel this is to imagine scaling the workload up. The launch tax is a fixed cost, say a few microseconds, paid no matter what. Now picture a slider for problem size. At two elements, the math takes a sliver of a nanosecond, so the runtime is essentially all tax, the GPU spends practically all its time on overhead, a total waste.

Push the slider up to a few hundred thousand elements and the math now takes long enough that the fixed tax is a thin sliver, a few percent. Push it to tens of millions of elements and the tax becomes invisible, lost in a flood of compute. This is the whole truth in one image: the GPU is not universally fast, it's fast only when the problem is big enough to bury its entry fee. Scale is the only real defense against a fixed cost.

THE SCALE EFFECT: FIXED COST VS WORK SIZETime(log)1 s1 ms1 µs1 ns21e51e7-1e8elements(~100k)(10M-100M)Work Size (more elements) →Total GPU TimeLaunch Overhead (~5µs)Actual Compute TimeKEY INSIGHTLaunch cost is afixed tax. Make thework big enoughto bury the tax.Scale is the onlyreal cure.

Dodging the Tax Directly

But when you can't just make the problem bigger, there are ways to dodge the tax directly.

1. Fusion: If you have three small pieces of work to do in a row, don't launch three separate times and pay the tax three times. Write one combined piece of work that does all three steps in sequence, paying the tax once. And there's a bonus that goes beyond saving launches. Without fusion, the first step writes its result out to slow main memory, and the second step reads it back, paying that long round trip in between. Fused into one, the first step's result stays right there in the fast on-chip registers, and the next step grabs it instantly without ever touching slow memory. So fusion saves both the launch tax and the data round trips between steps. The catch is that a bigger fused piece of work needs more registers per thread, and if it needs too many, you spill into slow memory and lose what you gained. There is a limit to how much you can usefully cram together.

AVOID THE FEE - (A) FUSIONBad: 3 launches = 3x the taxStep 1LaunchStep 2LaunchStep 3Launch3×taxGood: fuse them into one launchStep 1 → Step 2 → Step 3(one launch)1×taxBonus: it also skips the slow-memory round tripWithout fusionStep 1Step 2global mem(slow round trip)With fusionStep 1Step 2(stays on-chip, fast)Fuse to pay the tax once, and keep the data on-chip between steps.

2. Recording: In something like training a neural network, the same long chain of small operations runs over and over, thousands of times. Each time, the CPU walks the whole launch ceremony for every single operation, and the GPU often sits idle waiting for the CPU to dispatch the next one. The fix is to record the entire chain once into a single reusable object. The CPU does all the validation and setup one time during recording, building a fixed plan. Then every later run just replays that plan with a single command, bypassing nearly all the per-operation CPU overhead, and the gaps between operations collapse so they run tightly packed. The speedup doesn't come from faster math, the math is identical, it comes from eliminating the repeated CPU launch work. The catch is that the recorded plan freezes the shapes and addresses, so if your problem size keeps changing every run, the recording breaks and you have to rebuild it, which is expensive.

3. GPU Launching its Own Work: Instead of finishing a piece of work, returning to the CPU, and having the CPU launch the next piece, the GPU can spawn its own follow-up work from inside itself, useful when the next step depends on data only known at runtime. This removes the trip back to the CPU. But it does not make launching free. The GPU still has its own internal scheduling cost to start that work, and spawning many tiny pieces this way can actually be quite slow. It trades the CPU round trip for internal overhead, it doesn't erase the overhead.

THE DATA DANCE (YOU CAN'T SKIP IT)CPUSystem RAMGPUVRAM1. Copy In (H → D)2. Compute (on the GPU)kernel runs3. Copy Out (D → H)Copy in, compute, copy out - and these run SERIALLY.You cannot compute while the data is still copying.

↳ The General Takeaway

Launching GPU work is a formal cross-processor handoff that costs thousands of times more than a plain function call, a fixed entry fee paid before any math happens. That fee makes the GPU lose to a humble CPU core on small tasks, where the work is smaller than the tax.

The only true cure is scale, making the problem big enough to drown the fee. And where you can't, you fight it by fusing steps into one launch, recording repeated chains to replay cheaply, or letting the GPU launch its own follow-ups. The throughput of a GPU is enormous, but it is a heavy machine to start. You don't send it to deliver one letter. You send it when there is a mountain of mail.

THE BOTTOM LINETotal Time = Launch Overhead + Transfer In+ Compute Time + Transfer OutThe GPU wins only when:Compute Time >> (Launch + Transfer In + Transfer Out)High Math-to-Bytes OR Good Overlap OR Big Problem SizeIf compute can't dwarf the overhead, the work belongs on the CPU.