CUDA 13.1 Green Contexts Let Kernels Reserve SMs and Workqueues

NVIDIA has made green contexts available through the CUDA Runtime API starting with CUDA 13.1, giving applications explicit control over how GPU execution resources are divided within a single process. The feature, previously limited to the Driver API since CUDA 12.4, lets developers assign specific subsets of SMs and workqueues to a green context and target work to those resources directly.

How green contexts work

Green contexts address a problem that grows as GPU applications combine multiple independent components in one process: a latency-sensitive operator alongside a throughput-oriented background kernel, or a preprocessing stage sharing a GPU with model inference. Traditional CUDA contexts are heavyweight, incur hardware context-switch overhead, and were designed for an era when GPUs ran a single dominant workload.

With green contexts, an application creates a context for a chosen set of resources, then creates streams from that context. Work submitted to those streams is associated with the green context's resources. The Runtime API represents green contexts with the cudaExecutionContext_t type, and cudaGreenCtxCreate() returns a handle that can be passed directly to cudaExecutionCtxStreamCreate() instead of relying on implicit thread-local device state.

Beyond SM partitioning, green contexts can provision workqueue resources. In the traditional model, independent stream-ordered workloads may map to the same underlying workqueues, introducing unintended serialization even when execution resources are available. Explicit workqueue provisioning lets applications express expected concurrency and reduce false dependencies.

Measured latency impact

NVIDIA tested three modes with the same critical and bulk workload on a Blackwell GPU with 148 SMs. The critical kernel was a tiny integer workload; the bulk kernel was 20 back-to-back launches of a 4M-thread sqrtf loop designed to saturate the device.

In mode A, a green context partitioned 8 SMs for the critical kernel and 140 SMs for the bulk kernel, with the critical stream at high priority. Critical kernel latency was 0.007 ms. Mode B used the default context with a high-priority critical stream and no partition, yielding 0.140 ms. Mode C used the default context with equal priority on both streams, yielding 3.727 ms.

Stream priority alone delivered roughly a 27x improvement over equal-priority streams competing for the same SMs. But priority has a ceiling: an additional 20x was lost waiting for bulk blocks to drain, because stream priorities cannot preempt a block already executing on an SM. Green contexts skip that wait entirely because the critical SMs are dedicated to that workload, at the cost of fewer SMs for the bulk kernel.

Programming model and adoption

The code changes are minimal. Instead of cudaSetDevice() followed by cudaStreamCreate(), an application calls cudaGreenCtxCreate() with a resource descriptor, then cudaExecutionCtxStreamCreate() with the green context handle. The rest of the application continues using the same stream-based CUDA programming model.

Resource partitioning uses cudaDeviceGetDevResource() to query all SMs, cudaDevSmResourceSplit() to carve out a critical group aligned to the architecture's coscheduled alignment, and cudaDevResourceGenerateDesc() to pack each partition with its own workqueue-config resource. The example in NVIDIA's post also sets stream priorities on the green-context streams.

Green contexts are opt-in and additive. Existing applications can continue targeting the full device unchanged and adopt green contexts incrementally. NVIDIA suggests they are most useful when an application runs multiple independent workloads concurrently, needs more predictable performance rather than best-effort scheduling, wants explicit control over resource division, or sees degraded performance from limited workqueue control. The same skeleton can also be used to achieve desired concurrency, such as overlapping a communication kernel with a GEMM kernel so both make progress.