GPU Architecture · Execution & Operations
Many streams, many threads, many processes — all wanting one GPU. Who arbitrates, and when does work actually run together versus just take turns?
So far we've followed a single launch (Lesson 7) and a single process's async contract (Lesson 8). But your machine doesn't work that way in practice: a program has multiple streams, maybe multiple host threads, and the OS may have several processes all pointed at the same GPU. You asked how the system keeps track of all that — completions, failures, who's running. The answer has a surprising core: on a GPU, "running at the same time" is rarer than intent suggests, and one rule decides it — resources, not wishes.
Know the three levels of GPU sharing — streams inside a context, contexts across processes, and hardware partitions — and the single question that governs real overlap at every level: is there spare capacity, or are they taking turns?
Within a single process, independent work in different streams can run concurrently — the GPU's front end (Kepler-onward "Hyper-Q") offers up to 32 hardware work queues so streams don't create false dependencies on each other. But "can" is doing a lot of work in that sentence. Two kernels from the same context actually overlap only when both of these hold:1
That second condition is the reality check. A single well-sized kernel fills the GPU — enough blocks to cover every SM, or it maxes out registers/shared-memory per SM (recall occupancy, Lesson 2). It leaves nothing for a neighbour, so a second "ready" kernel simply waits its turn even though you launched it concurrently. Concurrent kernels mostly appear when the kernels are individually small (low occupancy, few blocks). Intent to overlap is free; the overlap itself is paid for in leftover capacity.
Now the case you specifically asked about — separate processes. Each process gets its own context (its own address space, allocations, streams — Lesson 7). And here is the load-bearing fact: work from different contexts cannot execute on the compute engine at the same time. By default the GPU time-slices — it gives one context a scheduled slice of the whole GPU, then context-switches to the next.2
Two processes share the GPU temporally: each gets the whole device for a slice, one at a time. They never truly run together — the driver switches contexts between slices.
So multi-process sharing is temporal, not spatial. As for "keeping track of it all": the driver holds per-context state — each context's own streams, its completion semaphores (Lesson 8), and its error/fault status. That's why a sticky fault poisons one context and leaves others alive: the tracking is per-context by construction. Completions and failures are bookkept the same way for every resident context; time-slicing just decides whose turn it is to make progress.
Two mechanisms break the "one context at a time" default when you genuinely need processes to share spatially:
MPS funnels both processes through one shared context, so their kernels run together on the same GPU — filling capacity a single process would waste.
| How it shares | Real concurrency? | Isolation | |
|---|---|---|---|
| Default | Temporal — time-slice per context | No — one context at a time | Strong (separate contexts) |
| MPS | Spatial — one shared context | Yes — cross-process kernels overlap | Weak (shared context; fault can cascade) |
| MIG | Spatial — hardware partition | Yes — fully independent instances | Strongest (dedicated hardware, QoS) |
One more sharing-adjacent idea, because it changes the submission model from Lesson 7. If a workload launches many small kernels in a repeating pattern, the per-launch CPU cost (that whole desugar → cudaLaunchKernel → pushbuffer → doorbell path, a few microseconds each) starts to dominate — the CPU becomes the bottleneck, submitting faster than it can keep up. CUDA graphs fix this: you capture the whole DAG of operations once, instantiate it once, then relaunch the entire graph with a single CPU call each iteration.5
cudaStreamBeginCapture(s, ...);
k1<<<...,0,s>>>(); k2<<<...,0,s>>>(); k3<<<...,0,s>>>(); // record, don't run
cudaStreamEndCapture(s, &graph);
cudaGraphInstantiate(&exec, graph, ...); // validate + build once
for (int i = 0; i < N; ++i) cudaGraphLaunch(exec, s); // ONE call replays the whole DAGInstead of re-submitting every op every iteration, the driver replays a pre-validated structure — and, seeing the whole workload up front, can schedule it better. It's the standard fix once profiling shows your kernels are tiny and the CPU can't feed them fast enough.
You're single-user on one GeForce GPU under WSL2, so you won't run MIG (unavailable) and probably not MPS. But the default behaviour is exactly what you'll see: launch two CUDA processes at once and they time-slice — each stutters, because they're taking turns on the whole GPU, not sharing it. And the within-context rule from §1 is the one you'll actually hit: adding streams won't speed up a kernel that already fills your 20 SMs. Graphs become relevant the moment you're launching thousands of small kernels in a loop.
From memory — one click locks each answer.
ConcurrencyYou launch two kernels into two separate non-default streams, but they still run one after the other. The most likely reason —
Different non-default streams permit concurrency, but it only happens if there are spare resources. A kernel that fills the SMs (blocks/registers/shared memory) leaves nothing for the second, so it waits — overlap mostly appears with small, low-occupancy kernels.
ProcessesBy default, two separate processes using the same GPU —
Work from different contexts can't run on the compute engine simultaneously, so the GPU time-slices: one context runs on the whole device, then it switches. True cross-process overlap needs MPS (shared context) or MIG (hardware partition).
GraphsCUDA graphs primarily help a workload that is —
Graphs amortize per-launch CPU overhead: capture the DAG once, then relaunch it with a single call. They help when kernels are small and numerous so submission cost dominates — not for one big compute- or bandwidth-bound kernel.
CUDA C++ Programming Guide — "Asynchronous Concurrent Execution" (concurrent kernels, streams, the resource conditions for overlap): read it here. For multi-process, the MPS documentation is the clearest statement of the time-slice default and what MPS changes. For graphs: Getting Started with CUDA Graphs.
One lesson left — and it's the one you came for. You now know how a launch runs (L7), how data and completion and errors flow (L8), and how the GPU is shared (L9). Lesson 10 closes the arc on your original goal: observability & debugging — what debug metadata travels in your binary, and exactly which data source each tool reads (nvidia-smi, Nsight Systems, Nsight Compute, compute-sanitizer, cuda-gdb), ending in a symptom→tool troubleshooting table. Ask anything here first, or say the word.
nvidia-cuda-mps-control, server nvidia-cuda-mps-server, client runtime in libcuda); a single shared context lets clients' kernels run concurrently; weaker fault isolation than separate processes (a fatal fault can affect co-resident clients). docs.nvidia.com/deploy/mps