Below the API

Kernel Launch and Host-Device Interaction

Intermediate Intermediate 1h Difficulty 3/5

Prerequisites 02, IV.08, II.09

1. What is it?#

Everything that happens at the CPU-GPU boundary: launching kernels, copying data, and synchronizing. This boundary costs microseconds, and at decode’s kernel granularity, microseconds matter.

Diagram — A launch is asynchronous#

sequenceDiagram
    participant P as Python / framework
    participant D as CUDA driver on CPU
    participant G as GPU
    P->>D: launch kernel
    D-->>P: returns immediately
    D->>G: enqueue on stream
    Note over P: CPU keeps going and queues more work
    G->>G: execute kernel
    P->>D: synchronize, .item() or .cpu()
    D-->>P: blocks until the stream drains
    Note over P,G: Time only after a sync, otherwise you timed the launch

2. Why does it exist as a topic?#

Because at batch 1, an LLM decode step is 200-2,000 kernels each doing ~10 µs of work, driven by a CPU issuing launches at ~5 µs each. The CPU can become the bottleneck for a $30,000 GPU.

Section IV.08 covered this from the framework’s perspective. This file covers the mechanics and the API-level tools.

3. Simple analogy#

Directing a construction crew by radio. Each instruction takes 10 seconds to transmit. If each task takes an hour, the radio time is irrelevant. If each task takes 15 seconds, you spend 40% of your day on the radio and the crew stands idle.

The fixes: batch instructions (“do these 50 things”), pre-record a work plan (CUDA graphs), or give the crew bigger tasks.

4. Tiny example#

import torch, time

def measure_launch_overhead():
    x = torch.randn(8, 8, device='cuda')      # tiny: work time ≈ 0
    torch.cuda.synchronize(); t0 = time.perf_counter()
    for _ in range(20000): y = x + 1
    torch.cuda.synchronize()
    return (time.perf_counter()-t0)/20000

print(f"launch+dispatch: {measure_launch_overhead()*1e6:.2f} us")

Typical: 5-15 µs on PyTorch (most of it Python dispatch, not the CUDA API itself). Pure C++ CUDA launch is ~3-5 µs.

Now the asynchrony:

x = torch.randn(8192, 8192, device='cuda')
t0 = time.perf_counter()
y = x @ x                                     # queues the kernel, returns immediately
t1 = time.perf_counter()
torch.cuda.synchronize()                      # NOW wait for it
t2 = time.perf_counter()
print(f"launch returned in {(t1-t0)*1e6:.1f} us; kernel took {(t2-t1)*1e3:.1f} ms")

launch returned in 42.3 us; kernel took 44.1 ms

The launch is 1000x cheaper than the work. This asynchrony is what lets the CPU stay ahead — until the kernels get small.

5. Technical explanation#

The launch path#

1. Application:      kernel<<<grid, block, shmem, stream>>>(args)
2. CUDA runtime:     validate, marshal arguments into a command packet
3. Driver:           write the packet into a pinned ring buffer in host memory
4. Doorbell:         a write to a memory-mapped GPU register
5. GPU front-end:    fetch the packet (DMA), decode
6. Work distributor: assign blocks to SMs as resources free up
7. Execution

Steps 1-4 are CPU work (~3-5 µs in C++, plus framework overhead). Steps 5-6 are GPU-side (~1-3 µs). They pipeline: while kernel N runs, the CPU can be launching N+1.

The failure mode: if step 1-4 takes longer than the kernel’s execution, the GPU drains its queue and idles.

Synchronization primitives#

torch.cuda.synchronize()              # block until ALL work on ALL streams is done
stream.synchronize()                  # block until this stream is done
event.synchronize()                   # block until this event is recorded
event.query()                         # non-blocking check
tensor.item(), .cpu(), print(tensor)  # IMPLICIT synchronization! ← the trap

Implicit synchronization is the most common performance bug in ML code. Any operation that reads a GPU value into Python forces the CPU to wait for the GPU to catch up, destroying the pipelining.

# BAD — synchronizes every step
for i in range(1000):
    loss = model(x)
    print(f"step {i}: {loss.item()}")     # .item() blocks!

# GOOD — accumulate on GPU, sync once
losses = []
for i in range(1000):
    losses.append(model(x))                # no sync
torch.cuda.synchronize()
print([l.item() for l in losses])

In inference: a per-token .item() to check a stopping condition costs ~0.5-1 ms per token. Keep stopping logic on the GPU, or check it every N tokens.

Memory transfers#

# Pageable (default): synchronous, ~10 GB/s
gpu_tensor = cpu_tensor.cuda()

# Pinned: asynchronous, ~25 GB/s
cpu_tensor = torch.empty(n, pin_memory=True)
gpu_tensor.copy_(cpu_tensor, non_blocking=True)

non_blocking=True only works from pinned memory. From pageable memory it’s silently synchronous — a common source of “why isn’t my overlap working?”

CUDA events for timing#

start = torch.cuda.Event(enable_timing=True)
end   = torch.cuda.Event(enable_timing=True)

start.record()
y = model(x)
end.record()
torch.cuda.synchronize()
print(f"{start.elapsed_time(end):.2f} ms")     # GPU-side timing, more accurate

Events measure GPU time directly, without including CPU launch overhead. Use events for kernel timing and wall-clock for end-to-end timing — and understand which question each answers.

Diagnosing launch-bound execution#

Symptom:  GPU utilization < 100% while a Python thread is at 100% CPU
Profile:  Nsight Systems timeline shows gaps between kernels
Metric:   sum(kernel durations) << wall clock time
Fix:      CUDA graphs (file 07), fusion, larger batch, move the loop out of Python

A quick check:

# If this ratio is much less than 1, you're launch/dispatch bound
total_kernel_time / wall_clock_time

PyTorch’s profiler gives you both directly.

6. Under the hood#

The GPU’s work distributor assigns blocks to SMs as slots free up. This means:

- Blocks from the SAME kernel can start at different times
- Blocks from DIFFERENT kernels on the same stream never overlap
  (unless the driver detects independence — it generally doesn't)
- Blocks from different STREAMS can overlap

That last point is why streams matter (file 06): a small kernel and a large kernel on separate streams can run concurrently, filling the GPU better than either alone.

7. Performance implications#

Scenario                              Launch overhead share
Prefill, S=2048, batch 8              < 1%
Decode, batch 128                     3-8%
Decode, batch 8                       15-30%
Decode, batch 1                       30-50%
Small model (< 1B), batch 1           50-70%

Small models at low batch are almost entirely launch-bound, which is why they benefit enormously from CUDA graphs and why running them on CPU is sometimes competitive.

8. Production implications#

  • Enable CUDA graphs. Biggest single fix (file 07).
  • Eliminate per-step .item(), .cpu(), and Python-side conditionals. Keep stopping criteria, sampling, and logits processing on the GPU.
  • Use pinned buffers for any recurring transfer, pooled and reused.
  • Provision enough CPU (Section II.01): 8-16 cores per GPU.
  • Profile with Nsight Systems to see gaps. This is the single most informative view of an inference server’s behavior.
  • Time with CUDA events in benchmarks, not time.time() without synchronization.

9. Common mistakes#

Timing without synchronize. Measures nothing.

Implicit sync via .item() in the hot loop. Adds ~1 ms per call.

non_blocking=True from pageable memory. Silently synchronous.

Assuming the GPU is the bottleneck without checking. At batch 1 it usually isn’t.

Allocating pinned memory per request. cudaHostAlloc is slow and serializing.

Printing tensors in a loop. Every print(tensor) is a synchronization.

10. Hands-on exercise#

A. Measure launch overhead. Run the microbenchmark in section 4. Compare PyTorch’s per-op overhead to a raw CUDA C++ launch if you can build one. Record in numbers.md.

B. Find the implicit syncs. Profile a model’s generation loop with Nsight Systems. Look for cudaStreamSynchronize or cudaMemcpyAsync D2H calls in the timeline. Trace each back to a Python line.

C. Compute the ratio. For a decode step, measure sum(kernel_time) and wall_clock. What fraction of wall clock is the GPU actually busy? Repeat at batch 1, 8, 64.

D. Pinned vs pageable. Measure transfer bandwidth from both, and verify that non_blocking=True only overlaps from pinned memory (use events to check overlap).

E. Fix a launch-bound loop. Take a small-model generation loop, measure it, then apply CUDA graphs. Report the speedup.

11. Interview questions#

  1. What happens when you launch a CUDA kernel? What does it cost?
  2. Why is CUDA execution asynchronous, and what breaks the asynchrony?
  3. Name four operations that cause implicit synchronization.
  4. When does non_blocking=True actually work?
  5. How would you determine whether a workload is launch-bound?
  6. Why are small models at batch 1 launch-bound, and what do you do about it?
  7. What’s the difference between timing with CUDA events and with wall clock?

12. Further reading#

  • [REFERENCE] CUDA C++ Programming Guide, “Asynchronous Concurrent Execution”
  • [REFERENCE] PyTorch CUDA semantics documentation
  • [REFERENCE] Nsight Systems user guide
  • Next: 06 — Streams and synchronization
esc

↑↓ navigate ↵ open