GPU（图形处理器） 应用程序越来越多地由多个独立组件组成，这些组件在同一进程内同时运行：一个延迟敏感的算子与一个面向吞吐量的后台内核并存；一个数据预处理阶段与模型推理并存；或者一个处理工作流的多个阶段共享同一块 GPU。

控制它们之间如何共享 GPU 资源仍然很困难。各组件可能以不可预测的方式相互干扰，而现有工具对资源进行分区的能力有限。

Green contexts 通过让应用程序显式选择 GPU 执行资源的子集，并将工作直接定向到这些资源，从而解决了这一问题。传统的 CUDA（NVIDIA 并行计算平台和编程模型） 上下文并非为此使用模型而设计。它们重量级、会引入硬件上下文切换开销，并且反映了早期时代的假设——当时 GPU 规模更小，应用程序通常作为单一主导工作负载运行。

Green contexts 自 NVIDIA CUDA 12.4 起已在 Driver API（应用程序编程接口） 中提供。从 CUDA 13.1 开始，它们也可通过 Runtime API 访问，允许应用程序在其进程内显式定义工作运行的位置以及执行资源如何划分。

Green contexts 的工作原理

Green contexts 的一个主要用途是 SM 分区。通过将特定的 SM 子集分配给一个 green context，应用程序可以将通过该 green context 提交的工作定向到这些 SM。这可以使多个工作负载在 GPU 上并发运行，而无需争用相同的计算单元。

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

Green contexts 的创建和销毁是轻量级的，创建或销毁一个 green context 不会隐式同步不相关的 GPU 工作。它们还提供了更显式的编程模型，应用程序将工作定向到选定的 green context，而不是仅依赖隐式/当前设备状态。

使用 green contexts 的显式编程模型

历史上，CUDA Runtime 应用程序通常使用 cudaSetDevice() 选择设备，创建流，并将工作提交到这些流。流的执行目标是在创建流时根据调用线程的当前设备或上下文推断的。

Green contexts 使这种定向更加显式。应用程序为选定的一组资源创建一个 green context，然后从该 green context 创建流。提交到这些流的工作与 green context 的资源相关联。

在 CUDA Runtime API 中，green contexts 由 cudaExecutionContext_t 类型表示，这是 CUDA 上下文的 Runtime 抽象。对于 green context 的使用，重要的一点是 cudaGreenCtxCreate() 返回一个句柄，可以直接传递给 cudaExecutionCtxStreamCreate() 等 API，而无需依赖隐式的线程本地设备或上下文状态。

对于流而言，这一转变直接体现在应用程序创建流的方式上。

注意：为简洁起见，以下代码片段省略了完整的错误检查。生产代码应检查所有 CUDA Runtime API 返回值，在内核启动后使用 cudaGetLastError() 或 cudaPeekAtLastError()，并检查 cudaStreamSynchronize() 等同步调用以捕获异步执行错误。

传统 Runtime 流程：流目标来自当前设备状态

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

Green-context 流程：流目标被显式选择

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

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

代码改动很小，但心智模型更清晰。应用程序不再依赖隐藏的线程本地设备状态，而是显式选择一个 green context 并为其创建流。应用程序的其余代码可以继续使用相同的基于流的 CUDA 编程模型。

希望面向整个设备的应用程序可以继续使用传统的 Runtime 模型。对于接受显式上下文句柄的 API，它们可以通过 cudaDeviceGetExecutionCtx() 获取设备级上下文。

示例

GPU 工作负载通常可能需要一个小型、延迟敏感的内核与一个面向吞吐量的大型工作线程共享设备。一个典型示例是分布式训练和推理中的通信/GEMM 重叠，或 NVIDIA Holoscan 等 AI 传感器处理平台中需要尽快启动的延迟敏感算子。

对此类关键工作进行优先级排序的标准工具之一是 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 上：

| 模式 | 关键内核延迟 |
| --- | --- |
| Green-ctx（关键使用 8 个 SM / 批量使用 140 个 SM） | 0.007 ms |
| 默认 ctx，高优先级关键 | 0.140 ms |
| 默认 ctx，相同优先级 | 3.727 ms |

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

仅使用流优先级相比不使用它就是一个巨大的胜利。比竞争相同 SM 的相同优先级流快约 27 倍。但优先级有上限。即使 GPU 调度器优先处理关键内核，仍有额外的 20 倍浪费在等待批量块排空上。Green contexts 完全跳过了这种等待，因为关键 SM 专用于该工作负载。代价是可用于批量内核执行的 SM 更少。

相同的骨架可用于实现所需的并发性而非延迟改进，例如将通信内核与 GEMM 内核重叠，使两者都能取得进展。高级用例通常同时使用 green context 分区和流优先级来达到所需的性能目标。

何时使用 green contexts

当你的应用程序满足以下条件时，它们最有用：

• 在同一 GPU 上并发运行多个独立工作负载。
• 需要更可预测的性能，而非尽力而为的流调度
• 希望对 GPU 执行资源的划分方式进行显式控制
• 由于对工作队列的控制有限而出现性能下降
• 正在向单进程内基于流水线或多组件的执行方式演进

入门

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

Green context 是可选启用且可叠加的。现有应用程序可以继续面向整个设备保持不变，并在需要对执行进行更精细控制的地方逐步采用 green context。

有关更多详细信息和示例，请参阅 CUDA Programming Guide 和 CUDA Runtime API 文档，并立即尝试 green context。

转载声明：本文为转载整理，原文版权归原作者及原站点所有。

原文来源：NVIDIA Technical Blog

原创地址：https://developer.nvidia.com/blog/control-how-your-gpu-shares-work-with-green-contexts/