GPU应用程序越来越多地由多个独立组件组成，这些组件在同一进程中同时运行：一个对延迟敏感的算子与一个面向吞吐量的后台内核并存；一个数据预处理阶段与模型推理并行；或者一个处理流程的多个阶段共享同一块GPU。
控制GPU资源在这些组件之间如何共享，仍然是一件困难的事。各组件之间可能出现难以预测的相互干扰，而现有工具对资源分区的支持能力也十分有限。
Green Contexts正是为解决这一问题而生，它允许应用程序显式选择一部分GPU执行资源，并将任务直接定向到这些资源上执行。传统的CUDA上下文并非为这种使用模式设计，它们较为笨重，存在硬件上下文切换开销，并且反映了早期GPU较小、应用程序通常以单一主导工作负载运行时代的设计假设。
Green Contexts自NVIDIA CUDA 12.4起已在驱动API中提供。从CUDA 13.1开始，它也可以通过运行时API（Runtime API）访问，使应用程序能够在其进程内，明确定义任务在何处运行，以及执行资源如何划分。
Green Contexts的工作原理
Green Contexts的一个主要用途是流多处理器（SM）分区。通过将特定的一组SM分配给某个Green Context，应用程序可以将通过该Green Context提交的任务定向到这些SM上。这使得多个工作负载能够在GPU上并发运行，而不会争夺相同的计算单元。
除了SM分区之外，Green Contexts还能配置工作队列资源。在传统模型中，独立的、按流排序的工作负载可能映射到相同的底层工作队列，即便拥有充足的执行资源，也会引入意外的串行化。通过显式配置工作队列，Green Contexts让应用程序能够表达其期望的并发程度，并减少虚假的依赖关系。
Green Contexts的创建和销毁都很轻量，创建或销毁一个Green Context不会隐式地同步无关的GPU任务。它们还提供了一种更为明确的编程模型，应用程序可以将任务定向到选定的Green Context，而不是仅依赖隐式的当前设备状态。
Green Contexts的显式编程模型
从历史上看，CUDA运行时应用程序通常通过cudaSetDevice()选择一个设备，创建流，然后向这些流提交任务。流的执行目标是在调用线程创建该流时，根据当前设备或上下文推断出来的。
Green Contexts让这种目标定位变得更加明确。应用程序先为所选的一组资源创建一个Green Context，然后从该Green Context创建流。提交到这些流中的任务将与该Green Context的资源相关联。
在CUDA运行时API中，Green Contexts由cudaExecutionContext_t类型表示，这是CUDA上下文的运行时抽象。对于Green Context的使用而言，关键点在于cudaGreenCtxCreate()会返回一个句柄，该句柄可以直接传递给诸如cudaExecutionCtxStreamCreate()这样的API，而不必依赖隐式的线程局部设备或上下文状态。
对于流而言，这种转变直接体现在应用程序创建流的方式上。
注：为简洁起见，以下代码片段省略了完整的错误检查。生产环境代码应检查所有CUDA运行时API的返回值，在内核启动后使用cudaGetLastError()或cudaPeekAtLastError()，并检查像cudaStreamSynchronize()这样的同步调用是否存在异步执行错误。
传统运行时流程：流的目标来自当前设备状态
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编程模型。
希望面向整个设备的应用程序可以继续使用传统的运行时模型。对于需要显式上下文句柄的API，可以通过cudaDeviceGetExecutionCtx()获取面向整个设备的上下文。
示例
GPU工作负载往往需要让一个延迟敏感的小型内核与一个面向吞吐量的大型工作负载共享设备。一个典型的例子是分布式训练和推理中的通信/GEMM重叠，或者像NVIDIA Holoscan这类AI传感器处理平台中需要尽快启动的延迟敏感算子。
优先处理这类关键任务的标准工具之一是CUDA流优先级。但遗憾的是，当批量内核已完全占满GPU的所有SM时，仅设置优先级并不足以保证高优先级内核立即执行。如果批量内核已经排队，而关键内核到达时，调度器会将关键内核的线程块分配给下一个空闲出来的SM。流优先级无法抢占已经在SM上执行的线程块，因此必须等待部分线程块执行完毕并释放资源。
我们来创建一个示例并查看结果。设置代码是我们需要编写的新内容（参见原文代码示例，包括查询设备SM资源、划分关键资源组、配置工作队列资源、生成资源描述符、创建Green Context及流等步骤）。
其余代码则应该是熟悉的样子（即在批量分区上饱和运行批量内核，在关键分区上启动关键内核，并通过事件记录测量其执行时间）。
流承载着分区信息，而内核只能看到它被允许运行的那些SM。
我们测试了三种模式，每种模式下的关键任务和批量任务工作负载相同：
模式A：Green Context分区，高优先级关键流
模式B：默认上下文，高优先级关键流（无分区）
模式C：默认上下文，两个流均为普通优先级（无流优先级设置，无分区）
关键内核是一个极小的整数计算任务。批量内核则是连续启动20次、每次400万线程规模的sqrtf循环，旨在使设备达到饱和状态。我们测量的是批量内核运行期间，关键内核的实际墙钟延迟。
在拥有148个SM的NVIDIA Blackwell GPU上测试结果显示：仅使用流优先级相较于不使用优先级就是巨大的提升，比优先级相同、争夺同一批SM的流快约27倍。但优先级存在上限——即便GPU调度器优先处理关键内核，仍有额外20倍的时间被浪费在等待批量任务线程块释放资源上。而Green Contexts完全跳过了这种等待，因为关键任务专用的SM是专属分配的。代价则是可用于批量内核执行的SM数量会减少。
同样的框架结构也可以用于实现预期的并发效果，而非仅仅改善延迟，例如让通信内核与GEMM内核重叠执行，使两者都能取得进展。在高级应用场景中，通常会同时使用Green Context分区和流优先级，以达成预期的性能目标。
何时应该使用Green Contexts
在以下情况下，Green Contexts最为有用：
应用程序需要在同一GPU上并发运行多个独立的工作负载；
需要更可预测的性能，而非依赖尽力而为式的流调度；
希望对GPU执行资源的划分方式拥有明确的控制权；
因工作队列控制能力有限而出现性能下降的情况；
正在朝着基于流水线或单进程内多组件执行的方向发展。
开始使用
使用CUDA 13.1或更新版本构建你的应用程序。
查询设备资源，并为你想要定向的资源分区创建一个Green Context。
为该Green Context创建流，并使用相同的基于流的CUDA编程模型提交任务。
对于需要显式上下文句柄且应面向整个设备的API，使用cudaDeviceGetExecutionCtx()。
Green Contexts是一种可选且附加的功能。现有应用程序可以继续保持面向整个设备不变，并在需要对执行进行更精细控制的地方逐步采用Green Contexts。
有关更多详情和示例，请参阅CUDA编程指南和CUDA运行时API文档，立即体验Green Contexts。
Q&A
Q1：Green Contexts是什么？它的主要作用是什么？
A：Green Contexts是NVIDIA CUDA提供的一种功能，允许应用程序显式选择一部分GPU执行资源并将任务定向到这些资源上，从而实现对GPU资源分区的精细控制，让多个工作负载能够在同一GPU上并发运行而互不干扰。
Q2：Green Contexts和CUDA流优先级有什么区别？
A：流优先级只能让调度器优先处理高优先级内核，但无法抢占已在SM上执行的线程块，仍需等待资源释放。而Green Contexts通过专属分配SM资源，让关键任务完全跳过等待，性能提升更显著。
Q3：如何开始使用Green Contexts？
A：需要使用CUDA 13.1或更新版本构建应用程序，查询设备资源并创建对应的Green Context，再从该Green Context创建流并提交任务，整体仍沿用基于流的CUDA编程模型。
