GPU 应用越来越多地由在单个进程中同时运行的多个独立组件组成：延迟敏感型运算符以及面向吞吐量的后台内核；数据预处理阶段以及模型推理；或共享单个 GPU 的处理工作流程的多个阶段。

仍然很难控制 GPU 资源在两者之间的共享方式。组件可能会造成难以预测的干扰，而现有工具对资源进行分区的能力有限。

绿色上下文通过允许应用显式选择 GPU 执行资源的子集，并将工作直接定位到这些资源来解决此问题。传统的 CUDA 上下文并非为此使用模型而设计。它们是重量级的，会产生硬件上下文交换开销，并且反映了早期时代的假设，当时 GPU 较小，应用程序通常作为单个主要工作负载运行。

自 NVIDIA CUDA 12.4 起，驱动 API 中便已提供绿色上下文。从 CUDA 13.1 开始，它们还可通过运行时 API 访问，从而允许应用程序在其进程中显式定义工作运行位置以及执行资源的分配方式。

绿色环境的工作原理

绿色上下文的一个主要用途是 SM 分区。通过将 SM 的特定子集分配给绿色环境，应用可以将通过该绿色环境提交的工作定向到这些 SM。这可以使多个工作负载在 GPU 上同时运行，而无需争夺相同的计算单元。

除了 SM 分区之外，绿色上下文还可以调配工作队列资源。在传统模型中，按流排序的独立工作负载可能会映射到相同的底层工作队列，即使有足够的执行资源，也会引入意外的序列化。通过显式置备工作队列，绿色上下文允许应用表达预期的并发性并减少虚假依赖项。

创建和销毁绿色上下文是轻量级的，创建或销毁绿色上下文并不会默认同步不相关的 GPU 工作。它们还提供了更明确的编程模型，其中应用将工作目标锁定在选定的绿色环境中，而不是仅依赖于隐式/ 当前设备状态。

具有绿色上下文的显式编程模型

过去，CUDA 运行时应用通常会选择具有 cudaSetDevice() 的设备，创建流并将工作提交给这些流。流的执行目标是从创建流时调用线程的当前设备或上下文中推断出的。

绿色上下文使目标更加明确。应用为所选资源集创建绿色上下文，然后根据该绿色上下文创建流。提交到这些流的工作与绿色环境的资源相关联。

在 CUDA 运行时 API 中，绿色上下文使用 cudaExecutionContext_t 类型 ( CUDA 上下文的运行时抽象) 表示。对于绿色上下文的使用，重要的一点是 cudaGreenCtxCreate() 返回可直接传递给 cudaExecutionCtxStreamCreate() 等 API 的句柄，而不是依赖隐式线程本地设备或上下文状态。

对于流，这种转变直接反映在应用创建流的方式上。

注意：为简洁起见，以下代码段省略了完整的错误检查。产品级代码应检查所有 CUDA 运行时 API 的返回值，在核函数启动后使用 cudaGetLastError() 或 cudaPeekAtLastError()，并检查同步调用 (例如 cudaStreamSynchronize()) 是否存在异步执行错误。

传统的运行时流：流目标来自当前设备状态

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

绿色上下文流：显式选择流目标

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

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

代码更改极少，但思维模型更清晰。应用显式选择绿色上下文并为其创建流，而不是依赖隐藏的线程本地设备状态。应用程序代码的其余部分可以继续使用相同的基于流的 CUDA 编程模型。

希望以整个设备为目标的应用可以继续使用传统的运行时模型。对于采用显式上下文句柄的 API，它们可以使用 cudaDeviceGetExecutionCtx() 获取设备范围的上下文。

示例

GPU 工作负载通常需要一个对延迟敏感的小型内核，以便与更大的吞吐量导向型工作者共享设备。一个典型示例是分布式训练和推理中的通信/ GEMM 重叠，或需要尽快启动的 AI 传感器处理平台 (如 NVIDIA Holoscan) 中对延迟敏感的运算符。

用于确定此类关键工作优先级的标准工具之一是 CUDA 流优先级。遗憾的是，当批量内核完全占用 GPU 的所有 SM 时，设置优先级不足以保证立即执行优先级更高的内核。如果批量核函数排成队列且关键核函数到达，调度程序会将关键核函数线程块传递给释放该函数的下一个 SM。流优先级无法抢占已在 SM 上执行的块，因此需要等待一些块耗尽。

我们来创建一个示例并查看一些结果。设置代码是我们需要编写的新部分：

// 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 

代码的其余部分看起来应该很熟悉：

// 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);

流承载分区，而内核只能看到允许其运行的 SM。

我们测试三种模式，每种模式的关键+ 批量工作负载相同：

• 模式 A：Green-ctx 分区，高优先级关键流
• 模式 B：默认上下文，高优先级关键流 (无分区)
• 模式 C：默认上下文，两个流的正常优先级 (无流优先级；无分区)

关键内核是一个微小的整数工作负载。批量内核是 20 个连续启动的 4M 线程 sqrtf 循环，旨在使设备饱和。我们会测量关键内核在批量运行时的时钟延迟。

在搭载 148 个 SM 的 NVIDIA Blackwell GPU 上：

表 1. 按执行模式划分的关键内核延迟

与不使用流优先级相比，单单流优先级本身就能带来巨大的胜利。与争夺相同 SM 的同等优先流相比，速度大约快 27 倍。但优先级有上限。即使 GPU 调度程序优先考虑关键内核，也会额外浪费 20 倍的时间等待大量块耗尽。绿色上下文完全跳过这种等待，因为关键 SM 专用于该工作负载。需要权衡一下，可用于批量核函数执行的 SM 会更少。

相同的骨架可用于实现所需的并发性，而非延迟改进，例如将通信核函数与 GEMM 核函数重叠，以便两者都能取得进展。高级用例通常同时使用绿色上下文分区和流优先级来实现所需的性能目标。

何时使用绿色上下文

它们在您的应用中非常有用：

• 在同一 GPU 上同时运行多个独立的工作负载。
• 需要更可预测的性能，而不是尽力而为的流调度
• 希望明确控制 GPU 执行资源的分配方式
• 发现由于对工作队列的控制有限而导致性能下降
• 正在向基于工作流或在单个流程中执行多组件的方向发展

开始使用

• 使用 CUDA 13.1 或更高版本构建您的应用。
• 查询设备资源，并为要定位的资源分区创建绿色上下文。
• 为该绿色环境创建流，并使用相同的基于流的 CUDA 编程模型提交工作。
• 对于采用显式上下文句柄且应面向整个设备的 API，请使用 cudaDeviceGetExecutionCtx()。

绿色背景是可选的，并且是叠加的。现有应用可以继续以整个设备为目标，保持不变，并在需要更精细地控制执行的地方逐步采用绿色环境。

有关更多详细信息和示例，请参阅CUDA 编程指南和CUDA 运行时 API文档，并立即试用绿色上下文。