← Chapter 26: Why the GPU context switch is free but CPU's is slow Chapter 28: Why NVIDIA made CUDA asynchronous →
date 15 / jul / 26
mtwtfss
SCALE & PERFORMANCE — Chapter 27 index home

The step-by-step anatomy of a kernel launch

*worked out from first principles*

You write one line to launch work on the GPU, and the instant you do, your main program moves on to the very next line. It does not wait. The GPU has not even started yet, and your CPU is already off doing something else. This single fact, that launching GPU work is fire-and-forget, is the source of a classic bug that makes beginners lose hours, and understanding it is what separates code that accidentally works from code that works on purpose.

Start with what actually happens when you launch. It feels like calling a function, where you call, it runs, it returns, and only then do you continue. GPU launches are nothing like that.

↳ When you launch, your CPU does not run the GPU work itself. It just writes a little note describing what to do, drops that note into a queue, and immediately walks away to the next line of your program. The note says which work to run, how many threads, and where the data is. That's it. The CPU is a manager scribbling an order onto a ticket and moving on, not the cook.

NOT A FUNCTION CALLNormal function call (CPU)result = add(a, b);add() runsreturnsNext line runsCall → Run → Return → ContinueGPU kernel launch (NOT a call)kernel<<<grid, block>>>(...)Write a ticketinto a queuetCPU immediatelycontinues to next lineLaunch → Enqueue → Continue (no wait)A launch drops a ticket in a queue - the CPU walks straight on, no waiting.

A relay of handoffs

And that note travels through a relay of handoffs before any math happens, no single part in charge of the whole trip. The CPU writes the ticket into a queue. The GPU's front-end eventually picks it up. A hardware distributor then hands whole blocks of work to whichever processing units have room. And inside each unit, schedulers feed the actual bundles of 32 threads into the cores.

Each stage passes to the next independently, which is exactly why it's a relay and not a single call. Nobody escorts the work end to end.

CPU Queue → GPU Front-end → Hardware Distributor → Scheduler → Cores

THE RELAY: HOW WORK GETS TO THE CORESCPU (Host)CPUWrites a ticket(command)into a queueGPU Front-EndFEEventuallypicks upthe ticketWork DistributorHands wholeblocks tofree SMsSM (Proc. Unit)SMScheduler runswarps (32threads)CoresMathfinallyhappens!Many independent handoffs. No one waits for the end-to-end trip.

The Stale-Read Bug

Now the consequence that bites everyone, and it follows directly from the CPU not waiting. Suppose you launch some work that fills an array on the GPU, and on the very next line you try to read the result. What do you get? Garbage. Because the GPU is still busy computing when your CPU races ahead and reads the memory.

The CPU is reading a result that does not exist yet. Two separate processors, one sprinting ahead of the other, and the reader gets there before the writer is done.

THE CLASSIC BUG: READING TOO EARLY// fill array on GPUkernel<<<...>>>(d_out);// immediately read resultprocess(d_out);// BUG!CPU Timelinelaunch kernelprocess(d_out)GPU Timelinekernel running...doneThe CPU reads before the GPU writes.You get GARBAGE (stale data).

The Hidden Checkpoint

So here's a puzzle. If that's true, why does so much beginner code seem to work anyway, reading results right after a launch without any obvious waiting?

Because something is quietly saving them without their knowledge. The moment you copy the results back from the GPU to the CPU, that copy operation refuses to begin until all the previously launched GPU work in its line has actually finished. So the copy itself acts as a hidden checkpoint.

WHY BEGINNERS THINK IT WORKSTypical beginner codekernel<<<...>>>(d_out);cudaMemcpy(h_out, d_out,size, ...DeviceToHost);process(h_out);What really happenscudaMemcpy (D→H) is synchronouswith respect to all prior work.kernelcudaMemcpy(D→H)CPU usesresultThe copy blocks until theGPU is done. It hides the bug?The copy-back is your accidental synchronization - it quietly hides the async bug.

Beginners never write an explicit wait, they just copy the data back, and the copy silently blocks until the GPU is done. The bug is there, but the copy hides it, which is why so many people never even learn that launches are asynchronous. Their code works by accident, propped up by a synchronization they didn't know was happening.

Waiting on Purpose

So how do you wait on purpose, when you do need to be sure the GPU has finished? There is a blunt instrument, a call that simply stops the CPU cold until every piece of GPU work issued so far has completed. It's the wait-for-everything hammer.

And there are lighter, more surgical versions that wait only for one specific line of work or one specific marker, which matter when you have several independent streams of GPU work running and you don't want to stall all of them just to check on one. The point is, waiting on the GPU is something you must ask for explicitly. It does not happen for free, and assuming it does is the bug.

HOW TO WAIT ON PURPOSEA) Heavy hammerwait for EVERYTHINGcudaDeviceSynchronize();W1W2W3W4...WnCPU waits until ALLprior GPU work is done.kgB) Targeted waitone thingcudaStreamSynchronize(s);or cudaEventSynchronize(e);S0:W1W2W3ES1:W4W5W6CPU waits onlyfor that stream.C) Ask laterquery, don't waitcudaStreamQuery(s);Returns immediately.Tells you if it's done.Not ready→ do elseReady→ use itWaiting is not free - block only when you actually need the result.

↳ The General Takeaway

Step back and the reason for all of this is one first principle we keep returning to: the CPU and the GPU are two separate computers joined by a bus. They don't share the same memory or the same flow of execution.

TWO SEPARATE COMPUTERSCPUCPUSystem RAMPCIe / BusGPUGPU Memory (VRAM)Send order→Do work→Report backThe CPU can't call the GPU - it sends a message across the bus, then awaits the reply.

So the CPU cannot call a GPU function the way it calls one of its own. It has to send a message across to a separate machine and, if it cares about the answer, explicitly wait for that machine to report back. The launch is the message, fired off and forgotten. The synchronization calls are how you await the reply. Thinking of it as a normal function call, where waiting is automatic, is precisely the mental model that produces the stale-read bug.

So the whole picture: launching GPU work is fire-and-forget. The CPU drops a ticket in a queue and immediately moves on while the work travels through a relay of independent handoffs before running. Because the CPU does not wait, reading results too early gives you garbage, the reader outrunning the writer. Beginners get rescued by the copy-back quietly blocking on their behalf, which hides the async nature until the day it doesn't. And the real fix is to synchronize on purpose, choosing the heavy wait-for-everything or a lighter targeted wait depending on what you actually need.

The GPU is not a subroutine you call and wait on.
It's a separate machine you send orders to and then, deliberately, ask when it's done.
Once you truly believe the CPU walked away, the whole model finally makes sense.