Control How Your GPU Shares Work with Green Contexts

GPU applications increasingly consist of multiple independent components running at the same time within a single process: a latency-sensitive operator alongside a throughput-oriented background kernel; a data preprocessing stage alongside model inference; or multiple stages of a processing workflow sharing a single GPU.

Controlling how GPU resources are shared between them remains difficult. Components can interfere unpredictably, and existing tools offer limited ability to partition resources.

Green contexts address this by letting applications explicitly select a subset of GPU execution resources and target work to those resources directly. Traditional CUDA contexts weren’t designed for this usage model. They are heavyweight, incur hardware context-switch overhead, and reflect assumptions from an earlier era, when GPUs were smaller and applications typically ran as a single dominant workload.

Green contexts have been available in the Driver API since NVIDIA CUDA 12.4. Starting with CUDA 13.1, they are also accessible through the Runtime API, allowing an application, within its process, to explicitly define where work runs and how execution resources are divided.

How green contexts work

One primary use of green contexts is SM partitioning. By assigning a specific subset of SMs to a green context, applications can target work submitted through that green context to those SMs. This can enable multiple workloads to run concurrently on the GPU without competing for the same compute units.

In addition to 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 sufficient execution resources are available. By provisioning workqueues explicitly, green contexts allow applications to express expected concurrency and reduce false dependencies.

Green contexts are lightweight to create and destroy, and creating or destroying one doesn’t implicitly synchronize unrelated GPU work. They also provide a more explicit programming model where applications target work to a chosen green context rather than relying only on implicit/current device state.

Explicit programming model with green contexts

Historically, CUDA Runtime applications typically selected a device with cudaSetDevice(), created streams, and submitted work to those streams. The stream’s execution target was inferred from the current device or context for the calling thread at the time the stream was created.

Green contexts make that targeting more explicit. An application creates a green context for a chosen set of resources, then creates streams from that green context. Work submitted to those streams is associated with the green context’s resources.

In the CUDA Runtime API, green contexts are represented with the cudaExecutionContext_t type, a Runtime abstraction for CUDA contexts. For green-context use, the important point is that cudaGreenCtxCreate() returns a handle that can be passed directly to APIs such as cudaExecutionCtxStreamCreate(), instead of relying on implicit thread-local device or context state.

For streams, this shift is reflected directly in how applications create them.

Note: For brevity, the following code snippets omit full error checking. Production code should check all CUDA Runtime API return values, use cudaGetLastError() or cudaPeekAtLastError() after kernel launches, and check synchronization calls such as cudaStreamSynchronize() for asynchronous execution errors.

Traditional Runtime flow: stream target comes from current device state

cudaSetDevice(device);
cudaStream_t s;
cudaStreamCreate(&s);
kernel<<<grid, block, 0, s>>>();

Green-context flow: stream target is selected explicitly

cudaExecutionContext_t greenCtx;
cudaGreenCtxCreate(&greenCtx, desc, device, 0);

cudaStream_t s;
cudaExecutionCtxStreamCreate(&s, greenCtx, 0, 0);
kernel<<<grid, block, 0, s>>>();

The code changes are minimal, but the mental model is clearer. Instead of relying on hidden thread-local device state, applications explicitly choose a green context and create streams for it. The rest of the application code can continue using the same stream-based CUDA programming model.

Applications that want to target the full device can continue using the traditional Runtime model. For APIs that take an explicit context handle, they can obtain the device-wide context with cudaDeviceGetExecutionCtx().

Example

GPU workloads may often require a small, latency-sensitive kernel to share a device with a larger throughput-oriented worker. A canonical example is communication/GEMM overlap in distributed training and inference, or latency-sensitive operators in AI sensor processing platforms like NVIDIA Holoscan that require starting as soon as possible.

One of the standard tools for prioritizing such critical work is CUDA stream priority. Unfortunately, setting priority isn’t enough to guarantee immediate execution of a higher priority kernel, when the bulk kernel fully occupies all of the GPU’s SMs. If bulk kernels are queued up and the critical kernel arrives, the scheduler hands critical kernel threads blocks to the next SM(s) that free(s) up. Stream priorities can’t preempt a block that’s already executing on an SM, so it needs to wait for some of the blocks to drain.

Let’s create a sample and look at some results. The setup code is the new piece we need to write:

// 1. Query all SMs on the device.
cudaDevResource all_sm {};
cudaDeviceGetDevResource(dev, &all_sm, cudaDevResourceTypeSm);

// 2. Carve out one critical group (arch's coscheduled alignment).
cudaDevSmResourceGroupParams crit_params {};
crit_params.smCount = all_sm.sm.smCoscheduledAlignment;

cudaDevResource critical_res{}, remaining_res{};
cudaDevSmResourceSplit(&critical_res, 1, &all_sm, &remaining_res, 0, &crit_params);

// 3. Pack each partition with its own workqueue-config resource so queue
//    pressure is isolated too, then generate the descriptor.
cudaDevResource wq {};
wq.type = cudaDevResourceTypeWorkqueueConfig;
wq.wqConfig.device = dev;
wq.wqConfig.sharingScope = cudaDevWorkqueueConfigScopeGreenCtxBalanced;
wq.wqConfig.wqConcurrencyLimit = 2;


cudaDevResource crit_pack[2] = { critical_res, wq };
cudaDevResourceDesc_t crit_desc{};
cudaDevResourceGenerateDesc(&crit_desc, crit_pack, 2);

cudaDevResource bulk_pack[2] = { remaining_res, wq };
cudaDevResourceDesc_t bulk_desc{};
cudaDevResourceGenerateDesc(&bulk_desc, bulk_pack, 2);

// 4. Create each green context, then a stream on it.
cudaExecutionContext_t crit_ctx{};
cudaGreenCtxCreate(&crit_ctx, crit_desc, dev, 0);

cudaStream_t crit_stream{};
cudaExecutionCtxStreamCreate(&crit_stream, crit_ctx, cudaStreamNonBlocking, prio_high);

cudaExecutionContext_t bulk_ctx{};
cudaGreenCtxCreate(&bulk_ctx, bulk_desc, dev, 0);

cudaStream_t bulk_stream{};
cudaExecutionCtxStreamCreate(&bulk_stream, bulk_ctx, cudaStreamNonBlocking, prio_low); // define prio_low 

While the rest of the code should look familiar:

// Saturate the bulk partition.
for (int i = 0; i < BULK_LAUNCHES; ++i) {
    bulk_kernel<<<bulk_grid, bulk_block, 0, bulk_stream>>>(
        d_bulk, N_BULK, BULK_ITERS);
}
// Launch the critical kernel on its own partition.
cudaEventRecord(t_start, crit_stream);
critical_kernel<<<crit_grid, crit_block, 0, crit_stream>>>(d_crit, N_CRIT);
cudaEventRecord(t_stop, crit_stream);
cudaStreamSynchronize(crit_stream);
float crit_ms = 0.0f;
cudaEventElapsedTime(&crit_ms, t_start, t_stop);

The stream carries the partition and the kernel sees only the SMs it’s allowed to run on.

We test three modes, same critical + bulk workload in each:

• Mode A: Green-ctx partition, high-priority critical stream
• Mode B: Default context, high-priority critical stream (no partition)
• Mode C: Default context, normal priority on both streams (no stream priorities; no partition)

The critical kernel is a tiny integer workload. The bulk kernel is 20 back-to-back launches of a 4M-thread sqrtf loop designed to saturate the device. We measure the wall-clock latency of the critical kernel while bulk is running.

On an NVIDIA Blackwell GPU with 148 SMs:

Table 1. Critical kernel latency by execution mode

Stream priority alone is a huge win over not using it. Roughly 27x faster than equal-priority streams competing for the same SMs. But priority has a ceiling. An additional 20x is wasted waiting for bulk blocks to drain, even with the GPU scheduler prioritizing the critical kernel. Green contexts skip that wait entirely because the critical SMs are dedicated to that workload. The tradeoff is that less SMs are available for the bulk kernel’s execution.

The same skeleton can be used to achieve desired concurrency rather than latency improvement, such as overlapping a communication kernel with a GEMM kernel so both make progress. Advanced use-cases often use both green context partitions and stream priorities to achieve desired performance targets.

When to use green contexts

They are most useful when your application:

• Runs multiple independent workloads concurrently on the same GPU.
• Needs more predictable performance rather than best-effort stream scheduling
• Wants explicit control over how GPU execution resources are divided
• Sees degraded performance due to limited control over workqueues
• Is evolving toward pipeline-based or multi-component execution within a single process

Getting started

• Build your application with CUDA 13.1 or newer.
• Query device resources and create a green context for the resource partition you want to target.
• Create streams for that green context and submit work using the same stream-based CUDA programming model.
• For APIs that take an explicit context handle and should target the full device, use cudaDeviceGetExecutionCtx().

Green contexts are opt-in and additive. Existing applications can continue to target the full device unchanged and adopt green contexts incrementally where finer control over execution is needed.

For additional details and examples, see the CUDA Programming Guide and CUDA Runtime API documentation, and try green contexts today.