GPU Architecture · Execution & Operations
The launch returned to your CPU thread immediately. So how did the inputs get there, how do you learn it finished, and why do errors surface late?
Lesson 7 ended on a cliffhanger: cudaLaunchKernel is asynchronous — it enqueues the work and hands control straight back to the CPU, usually before the kernel has even started. That one fact forces three questions, and all three have the same shape of answer: the CPU and GPU run decoupled, talking through ordered queues and completion signals rather than lockstep calls. This lesson is those three questions.
Internalize the asynchronous host↔device contract: bytes move by DMA (and pinned memory is why it's fast), completion arrives through streams and events (not by the call blocking), and errors are async and can be "sticky" (which is why a crash seems to happen on the "wrong" line).
cudaMemcpy ships data across the PCIe link, but the CPU doesn't do the copying — a dedicated DMA engine (a "copy engine") on the GPU reads and writes host RAM over PCIe while your CPU is free. There's a catch that explains a whole class of "why is my transfer slow" questions: a DMA engine addresses host memory by physical address and needs those pages to stay put. Ordinary host memory is pageable — the OS can relocate or evict it anytime — so the driver can't hand its address to the DMA engine directly. Instead it does a hidden two-step:1
cudaMallocHost)Page-locked ("pinned") memory is host memory the driver has locked in place, so the DMA engine can read it directly — skipping the staging "bounce buffer." That's why pinned transfers are meaningfully faster, and, more importantly, why they're the only ones that can go truly asynchronous: cudaMemcpyAsync on pageable memory silently falls back to blocking, because the driver still needs that serialized staging copy. Pinned memory is the price of admission for overlapping transfers with compute.1
cudaMallocHost/cudaHostAlloc (or pin an existing region with cudaHostRegister).
Every async operation you issue — a launch, an cudaMemcpyAsync — goes into a stream: an ordered, in-order queue of GPU work. Operations in one stream execute strictly in enqueue order; that's the ordering guarantee you build on. To learn when work is finished, you have three documented tools:2
cudaStreamSynchronize(stream) — block this CPU thread until everything queued in that stream has completed. cudaDeviceSynchronize() is the bigger hammer: wait for all streams.cudaStreamQuery(stream) — the non-blocking poll: cudaSuccess if drained, cudaErrorNotReady if not. Lets the CPU do other work meanwhile.cudaEventRecord(e, stream) enqueues the marker; cudaEventSynchronize(e) waits for the stream to reach it. And cudaEventElapsedTime(&ms, start, stop) gives you the interval between two events measured by GPU-recorded timestamps — which is the correct way to time a kernel, immune to the async launch gap and CPU jitter.3How does cudaStreamSynchronize know to return? The GPU, having executed your commands in order, writes a rising value to a known memory location (a semaphore/fence) as its last step — so the moment that value appears, everything before it is provably done. The waiting CPU thread learns of it one of two ways, and you can choose via cudaSetDeviceFlags: spin (cudaDeviceScheduleSpin — busy-poll that memory location, lowest latency, pins a core at 100%) or block (cudaDeviceScheduleBlockingSync — sleep until the GPU raises a PCIe interrupt that the driver services to wake you, ~0% CPU). The spin-vs-block choice and its CPU-usage signature are documented and directly observable; the interrupt/semaphore plumbing itself is pieced together from NVIDIA's open driver, not the CUDA manual. Watch a synchronizing program's CPU use flip between ~100% and ~0% as you toggle the flag — that's the two paths, made visible.4
One stream gives you order. Multiple streams give you overlap — the real payoff. Work in different streams has no ordering between them and may run concurrently, so while one stream's kernel computes, another's copy engine can be moving the next chunk of data. With pinned memory and separate streams, a naïvely serial pipeline collapses:
copy (H2D / D2H) kernel. Splitting the work into chunks across streams lets copies and compute overlap — the job finishes in roughly half the wall-clock. This is the whole reason streams and pinned memory exist.
Two things gate that picture in practice. First, copy engines: cudaDeviceProp::asyncEngineCount tells you how much overlap the hardware allows — 0 none, 1 overlap copy with compute, 2 also overlap the two copy directions at once. Read it from your deviceQuery; consumer parts often expose fewer engines than datacenter ones.5 Second, the default-stream trap: the legacy default stream (plain <<<>>> with no stream argument) implicitly synchronizes with all other streams — slip one operation onto it between two others and you silently serialize work you meant to overlap. For a genuine cross-stream dependency, express it explicitly with cudaStreamWaitEvent(streamB, event) — "B waits for the event A recorded" — instead of falling back to the default stream.2
Unified Memory (cudaMallocManaged) sidesteps explicit copies: one pointer valid on both sides, and pages fault in and migrate on demand (Pascal and newer). Convenient — but on WSL2, concurrent CPU/GPU access to managed memory is unsupported, and the extra VM-bus hop raises small-transfer overhead. On your setup, prefer explicit cudaMemcpy with pinned buffers when you're measuring, and treat managed memory as a convenience, not a performance path.6
Here's the payoff for troubleshooting, and the thing that confuses everyone once. Because a launch returns before the kernel runs, a fault inside the kernel hasn't happened yet when the launch call returns. So CUDA splits errors into two kinds:7
| Kind | When it surfaces | Examples |
|---|---|---|
| Launch / config (synchronous) | At the launch itself | too many threads/block, bad launch config, invalid stream |
| Execution (asynchronous) | Later — on the next sync or CUDA call | illegal address (700), misaligned access, device-side assert/trap |
You retrieve them with two functions that differ in one way: cudaGetLastError() returns the last error and clears it; cudaPeekAtLastError() returns it without clearing. The standard idiom brackets a launch — clear stale state, launch, check the config error, then sync to force and check the execution error:
cudaGetLastError(); // clear any stale error first
myKernel<<<grid, block>>>(args);
cudaError_t cfg = cudaGetLastError(); // LAUNCH/config error (synchronous)
cudaError_t run = cudaDeviceSynchronize(); // EXECUTION error (surfaces here)Now the sharp edge. Execution faults like an illegal address are sticky: cudaGetLastError can't clear them, because they've corrupted the entire context. After one, every subsequent CUDA call in the process — even an unrelated cudaMalloc — returns that same error. The context is dead; the only real recovery is to tear it down (cudaDeviceReset, situationally) or, reliably, restart the process.7
badKernel<<<g, b>>>(...); // reads a bad address deep inside
cudaDeviceSynchronize(); // → cudaErrorIllegalAddress (700)
cudaMalloc(&p, 16); // → STILL 700. context is poisonedWhen a CUDA error appears on a line that "can't" be wrong — a cudaMalloc, a second kernel — suspect an earlier async fault. The reported line is just where the poisoned context was next touched, not where the crash happened. To find the real culprit: check errors right after each launch with a sync, or run under compute-sanitizer (Lesson 6), which pins the fault to the exact kernel and address.
From memory — one click locks each answer.
MechanismPinned (page-locked) host memory speeds up transfers mainly because —
A DMA engine needs fixed physical addresses. Pageable memory can be relocated, so the driver first copies it to a hidden pinned staging buffer, then DMAs that. Pinning removes the extra copy — and is what lets cudaMemcpyAsync actually run async.
CompletionThe cleanest way to time a kernel accurately is —
The launch is async — a CPU timestamp after the launch line captures almost nothing, since the kernel hasn't run yet. Events are recorded in the stream and timed by GPU timestamps, so they measure actual on-GPU execution.
ErrorsA cudaMalloc suddenly returns cudaErrorIllegalAddress. The most likely cause is —
Illegal-address is a sticky execution error from a kernel. Once it corrupts the context, every later call — including an unrelated cudaMalloc — returns the same code. The malloc line is where you noticed it, not where it happened.
Mark Harris, "How to Optimize Data Transfers in CUDA C/C++" (pinned memory, the bounce buffer, async overlap) — read it here. Then the CUDA C++ Programming Guide, "Asynchronous Concurrent Execution" for streams, events, and overlap semantics.
You now know how one process talks to the GPU. But a real machine runs many streams, many threads, even many processes against one GPU at once — who arbitrates, and can they truly run together? And once it's all running, how do you actually see inside it to troubleshoot? That's Lesson 9 — multi-tenancy (streams, contexts, MPS, MIG, CUDA graphs) and observability (what metadata travels, and exactly what data each tool reads). It closes the arc on your original goal: troubleshooting. Ask anything here first, or say the word.
cudaStreamWaitEvent for cross-stream dependencies) and "Stream Synchronization Behavior" (the legacy default stream synchronizes with other blocking streams). docs.nvidia.comcudaEventRecord/cudaEventSynchronize/cudaEventQuery; cudaEventElapsedTime is derived from GPU-recorded timestamps and needs timing-enabled events. docs.nvidia.comcudaSetDeviceFlags with cudaDeviceScheduleSpin / cudaDeviceScheduleYield / cudaDeviceScheduleBlockingSync selects busy-poll vs interrupt-blocked waiting (documented, observable via CPU usage). The underlying semaphore-release + MSI-interrupt completion path is reconstructed from NVIDIA's open GPU kernel modules and Linux MSI documentation — not the CUDA programming guide. docs.nvidia.comcudaDeviceProp::asyncEngineCount (0 = no overlap; 1 = overlap copy with kernel; 2 = also bidirectional copy overlap). Read it from deviceQuery; it supersedes the deprecated deviceOverlap field. docs.nvidia.comcudaGetLastError returns and clears the last error; cudaPeekAtLastError returns without clearing; error functions may report errors from prior asynchronous launches. Sticky execution errors corrupt the context (every subsequent call returns the same error); recovery is context teardown / process restart. Bob Crovella, "CUDA Debugging" (NVIDIA/OLCF training). docs.nvidia.com