Lesson 9 · Execution & operations arc

GPU Architecture · Execution & Operations

Sharing the GPU: concurrency & multi-tenancy

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.

The one win

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?

1 · Inside one context: concurrency needs room

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.

This is why "just use more streams" rarely speeds up a program whose kernels are already large: they were saturating the machine anyway. Streams help when you have many small or heterogeneous pieces (e.g. a copy that can ride alongside a compute-light kernel).

2 · Across processes: taking turns by default

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

Default · time-slice
proc A
proc B
proc A
proc B

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.

3 · Going further: MPS and MIG

Two mechanisms break the "one context at a time" default when you genuinely need processes to share spatially:

MPS · shared context
proc A
 
proc B

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 sharesReal concurrency?Isolation
DefaultTemporal — time-slice per contextNo — one context at a timeStrong (separate contexts)
MPSSpatial — one shared contextYes — cross-process kernels overlapWeak (shared context; fault can cascade)
MIGSpatial — hardware partitionYes — fully independent instancesStrongest (dedicated hardware, QoS)

4 · CUDA graphs: when the bottleneck is you

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 DAG

Instead 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.

On your box

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.

Check yourself

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 —

ProcessesBy default, two separate processes using the same GPU —

GraphsCUDA graphs primarily help a workload that is —

Primary source (≈15 min)

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.

Notes & citations

  1. CUDA C++ Programming Guide, "Concurrent Kernel Execution" — kernels from one context may execute concurrently given separate (non-default) streams and available resources; Hyper-Q (Kepler+) provides up to 32 hardware work queues. Resource availability is the practical limiter. docs.nvidia.com
  2. NVIDIA MPS documentation — "work launched to the compute engine from work queues belonging to different CUDA contexts cannot execute concurrently"; without MPS the GPU time-slices, assigning each process a serially-scheduled slice of the whole GPU. docs.nvidia.com/deploy/mps
  3. NVIDIA MPS documentation — architecture (control daemon 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
  4. NVIDIA Multi-Instance GPU User Guide — partitions a GPU into up to 7 instances with dedicated SMs, memory, and isolated memory-system paths; introduced on A100 and available on datacenter parts (A100/A30/H100/…), not GeForce. docs.nvidia.com/datacenter/tesla/mig-user-guide
  5. NVIDIA Developer Blog, "Getting Started with CUDA Graphs" — define → instantiate → launch; capturing a repeated sequence of small launches into one graph amortizes per-launch CPU submission overhead (illustrative figures are from a specific example, not universal constants). developer.nvidia.com
← Lesson 8 ⌂ Home Lesson 10 →