Green Contexts

重點總覽

Green Context 是什麼

green context (GC) 是一個 輕量 context，從建立那一刻起就與一組特定 GPU 資源綁定。使用者可在建立時 partition GPU 資源（目前是 streaming multiprocessors, SMs 與 work queues, WQs），使得 targeting 該 GC 的 GPU 工作 只能使用其 provisioned 的 SM 與 work queue。如此可減少或更好地控制因共用資源造成的干擾。一個應用程式可擁有多個 green context。

• 使用 green context 不需要任何 GPU code（kernel）變更，只需小幅 host 端改動（建立 green context、為其建立 stream）。
• 典型情境：確保某些 SM 永遠可供 latency-sensitive kernel 立即啟動（在無其他限制下）；或快速測試「使用較少 SM」的效果而不改 kernel。
• green context 最早經由 CUDA Driver API 提供；CUDA 13.1 起 經由 execution context (EC) 抽象暴露在 CUDA runtime。

Motivation / 何時使用

啟動 CUDA kernel 時，使用者 無法直接控制 kernel 會在幾顆 SM 上執行，只能透過改變 launch geometry（或任何影響每 SM 最大 active block 數的因素）間接影響。當多個 kernel 在 GPU 上並行（不同 stream 或 CUDA graph 的一部分），它們也可能 爭用相同 SM 資源。

green context 透過 partition SM 資源 解決此問題：某 GC 只能用建立時 provisioned 的特定 SM。

無 Green Context（kernel A 先佔滿全部 SM）
SM% │████████████████ A ███████████──┐
    │                    B 被延遲 →   └─B 等 A ramp down 才能開始
    └───────────────────────────────────────► time

有 Green Context（A 限 80% SM，B 保留專屬 SM）
SM% │████████ A ████████████████████
    │░░░░ B 立即開始（有專屬 SM）░░░░
    └───────────────────────────────────────► time

• 上圖對應原文 Figure 45：A 先啟動佔滿 SM，latency-sensitive 的 B 後啟動時無 SM 可用，必須等 A 的 thread block 結束。
• 用 GC 後，A 只能用 GC-A 的 SM（不論 launch config），B 被啟動時 保證有可用 SM 立即開始（除非有其他資源限制）。A 的 duration 可能因此變長，但 B 不再被「無 SM」延遲。
• 每個 GC 能用幾顆 SM 應由使用者在建立時 逐案決定。

Work Queues 的角色

SM 是可 provisioned 的一種資源；另一種是 work queues (WQs)。把 work queue 想成一個黑箱資源抽象，它（連同其他因素）會影響 GPU 工作的並行度。

• 若獨立的 GPU 工作（如不同 stream 上的 kernel）映射到 同一個 work queue，可能引入 false dependence，導致序列化執行。
• 即使 SM 可用，若 B 與 A 映射到同一 WQ，B 仍可能必須等 A 完整結束。
• 使用者可透過 CUDA_DEVICE_MAX_CONNECTIONS 環境變數影響 GPU 上 work queue 數量上限。
• GC 讓使用者用「預期的並行 stream-ordered workload 數」表達最大並行度；driver 以此為 hint，盡量避免不同 EC 的工作用到同一 WQ。

GC vs MIG vs MPS

• MIG 無法 解決「同一應用內 B 被佔滿 SM 的 A 延遲」的問題（單一 instance 內仍會發生）；解法是在 MIG 上 搭配 GC。
• MPS 的 active thread percentage 表示「不超過 x% 的 SM」但那 N 顆可以是任意且隨時間變；GC 的 N 顆是 固定特定 的 N 顆。
• 透過 cuCtxCreate driver API + execution affinity 也能做 programmatic SM partition，但同樣是任意 N 顆、會變動。
• 建立 GC 遠比 MPS context 輕量，因許多底層結構由 primary context 擁有並共享。

Ease of Use（最小改動）

原本只能靠 launch geometry 間接影響 SM 用量：

int gpu_device_index = 0;  // GPU ordinal
cudaSetDevice(gpu_device_index);
cudaStream_t strm1, strm2;
cudaStreamCreateWithFlags(&strm1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&strm2, cudaStreamNonBlocking);
// 無法控制各 stream 上 kernel 能用幾顆 SM
code_that_launches_kernels_on_streams(strm1, strm2);

CUDA 13.1 起，只要加幾行 host code、不改 kernel，即可控制每個 stream 上 kernel 可用的 SM：

// 1) 取得全部 SM 資源
cudaDevResource initial_GPU_SM_resources{};
cudaDeviceGetDevResource(gpu_device_index, &initial_GPU_SM_resources,
                         cudaDevResourceTypeSm);
// 2) split：一組 16 SM、一組 8 SM
cudaDevSmResource result[2]{{}, {}};
cudaDevSmResourceGroupParams group_params[2] = {
  {.smCount=16, .coscheduledSmCount=0, .preferredCoscheduledSmCount=0, .flags=0},
  {.smCount=8,  .coscheduledSmCount=0, .preferredCoscheduledSmCount=0, .flags=0}};
cudaDevSmResourceSplit(&result[0], 2, &initial_GPU_SM_resources, nullptr, 0,
                       &group_params[0]);
// 3) 各自產生 descriptor
cudaDevResourceDesc_t resource_desc1{}, resource_desc2{};
cudaDevResourceGenerateDesc(&resource_desc1, &result[0], 1);
cudaDevResourceGenerateDesc(&resource_desc2, &result[1], 1);
// 4) 建立兩個 green context
cudaExecutionContext_t my_green_ctx1{}, my_green_ctx2{};
cudaGreenCtxCreate(&my_green_ctx1, resource_desc1, gpu_device_index, 0);
cudaGreenCtxCreate(&my_green_ctx2, resource_desc2, gpu_device_index, 0);
// 唯一需改的呼叫端：改用 EC 版 stream 建立 API
cudaStream_t strm1, strm2;
cudaExecutionCtxStreamCreate(&strm1, my_green_ctx1, cudaStreamDefault, 0);
cudaExecutionCtxStreamCreate(&strm2, my_green_ctx2, cudaStreamDefault, 0);
// 其餘不變：strm1 上 kernel 最多 16 SM、strm2 上最多 8 SM
code_that_launches_kernels_on_streams(strm1, strm2);

唯一的呼叫端差異是 用 cudaExecutionCtxStreamCreate 取代 cudaStreamCreateWithFlags，函式內部與 kernel 完全不動。

Device Resource 與 Resource Descriptor

green context 的核心是綁定特定 device 的 device resource (cudaDevResource)；資源可被組合並封裝進 descriptor (cudaDevResourceDesc_t)。GC 只能存取建立時 descriptor 內封裝的資源。

struct {
  enum cudaDevResourceType type;
  union {
    struct cudaDevSmResource              sm;        // SM 型
    struct cudaDevWorkqueueConfigResource wqConfig;  // WQ config 型
    struct cudaDevWorkqueueResource       wq;        // pre-existing WQ 型
  };
};

• 有效 type：cudaDevResourceTypeSm、cudaDevResourceTypeWorkqueueConfig、cudaDevResourceTypeWorkqueue；cudaDevResourceTypeInvalid 代表無效。
• 一個 GPU device 預設具備全部三型：涵蓋全部 SM 的 SM 資源、涵蓋全部 WQ 的 workqueue config 資源、及其對應的 workqueue 資源；皆可由 cudaDeviceGetDevResource 取得。
• 查詢某 EC 或 stream 是否關聯某型資源：cudaExecutionCtxGetDevResource（EC，可關聯多型）、cudaStreamGetDevResource（stream，只能關聯 SM 型）。
• 建議所有 device resource struct 都 zero-initialize。

• sharingScope 兩值：cudaDevWorkqueueConfigScopeDeviceCtx（預設，所有 WQ 資源跨所有 context 共享）與 cudaDevWorkqueueConfigScopeGreenCtxBalanced（driver 盡量在各 GC 間用 不重疊 的 WQ，以 wqConcurrencyLimit 為 hint）。

建立 Green Context：四步驟

Step 1            Step 2                    Step 3                Step 4
取得資源    →    split SM 資源        →    產生 descriptor   →   建立 GC
GetDevResource   SplitByCount / Split      GenerateDesc          GreenCtxCreate
(device/EC/      (+ 選配 WQ config)        (封裝一或多 resource) (provision 資源)
 stream)
                                                                     │
                            建立屬於 GC 的 stream ◄───────────────────┘
                            ExecutionCtxStreamCreate → <<< >>> 只用該 GC 資源

Step 1：取得可用 GPU 資源

三個可能起點：device、execution context、stream。

// device:
cudaError_t cudaDeviceGetDevResource(int device, cudaDevResource* resource,
                                     cudaDevResourceType type);
// execution context:
cudaError_t cudaExecutionCtxGetDevResource(cudaExecutionContext_t ctx,
                                           cudaDevResource* resource,
                                           cudaDevResourceType type);
// stream（只支援 SM 型）:
cudaError_t cudaStreamGetDevResource(cudaStream_t hStream,
                                     cudaDevResource* resource,
                                     cudaDevResourceType type);

• 三者都接受所有有效 type，唯一例外 是 cudaStreamGetDevResource 只支援 SM 型。
• 通常起點是 GPU device。成功取得 SM 資源後可讀 sm.smCount、sm.minSmPartitionSize、sm.smCoscheduledAlignment。
• 取得 WQ config 資源時，若起點是 device，wqConcurrencyLimit 會等於 CUDA_DEVICE_MAX_CONNECTIONS 的值或其預設。

Step 2：partition SM 資源

把可用 SM 資源 靜態 split 成一或多個 partition（可能留下一個 remaining partition）。兩個 split API 都僅適用 SM 型資源：

SplitByCount：請求把 input 切成 *nbGroups 組、每組 minCount SM。實際結果的 *nbGroups 可能變小（≤ 請求值）、實際 N 可能變大（≥ minCount），因架構特定的 granularity/alignment 要求。CC 9.0 範例：每 partition 最少 8 SM、SM 數須為 8 的倍數（useFlags=0）。

unsigned int min_SM_count = 8;
unsigned int actual_split_groups = 5;  // 可能被更新
cudaDevResource actual_split_result[5] = {{}, {}, {}, {}, {}};
cudaDevResource remaining_partition = {};
cudaDevSmResourceSplitByCount(&actual_split_result[0], &actual_split_groups,
                              &avail_resources, &remaining_partition,
                              0 /*useFlags*/, min_SM_count);

• result=nullptr 可用來 查詢 會建立幾組；remaining=nullptr 表示不在乎 leftover。
• remaining（leftover）partition 不具備 與同質 group 相同的功能/效能保證。
• useFlags 預設 0，亦支援 cudaDevSmResourceSplitIgnoreSmCoscheduling、cudaDevSmResourceSplitMaxPotentialClusterSize。
• 任何 split 出的 cudaDevResource 無法直接再 split，必須先以它建 descriptor 與 GC（即先做 Step 3、4）。

Split（異質）：單次呼叫產生 non-overlapping 異質 partition，放入 result 陣列，各組 SM 數可不同但 絕不為 0。每組以 cudaDevSmResourceGroupParams 指定需求：

• discovery mode：把某組 smCount 設 0，成功後該 groupParams.smCount 會被填入有效非 0 值（actual smCount）；若 result 非 null，對應 group 的 sm.smCount 也設為同值。
• groupParams 各 entry 由左（index 0）到右評估，順序重要（順序不同可能導致後面的組沒 SM 可用）。
• coscheduledSmCount／preferredCoscheduledSmCount 設 0 表示用預設值，兩者預設皆等於該 device（透過 cudaDeviceGetDevResource 取得的）SM 資源之 smCoscheduledAlignment。
• 回傳值依 result：result != nullptr 時只有成功建立 nbGroups 個有效組才回 cudaSuccess，否則回錯（不同錯誤可能同碼如 CUDA_ERROR_INVALID_RESOURCE_CONFIGURATION，開發時用 CUDA_LOG_FILE 取得更詳細描述）；result == nullptr 是 dry-run，即使某組 smCount 為 0 也可能回 cudaSuccess，適合探索。
• remainder（若存在）對 SM 數與 coschedule 無任何約束，由使用者自行探索。

Step 2（續）：加入 workqueue 資源

若也要指定 WQ 資源，須 顯式 建立（無類似 split 的 API）：

cudaDevResource split_result[2] = {{}, {}};
// split_result[0] 由 split API 填（nbGroups=1）；最後一個是 WQ 資源
split_result[1].type = cudaDevResourceTypeWorkqueueConfig;
split_result[1].wqConfig.device = 0;                                     // device ordinal
split_result[1].wqConfig.sharingScope = cudaDevWorkqueueConfigScopeGreenCtxBalanced;
split_result[1].wqConfig.wqConcurrencyLimit = 4;                         // hint：最多 4 並行

wqConcurrencyLimit = 4 提示 driver 預期最多 4 個並行 stream-ordered workload；driver 盡量在可能時尊重此 hint 分配 work queue。

Step 3：建立 Resource Descriptor

cudaError_t cudaDevResourceGenerateDesc(cudaDevResourceDesc_t* phDesc,
                                        cudaDevResource* resources,
                                        unsigned int nbResources);

// 把 actual_split_result[2] 到 [4] 共 3 個資源封裝進一個 descriptor
cudaDevResourceDesc_t resource_desc;
cudaDevResourceGenerateDesc(&resource_desc, &actual_split_result[2], 3);

成功條件：

• 全部 nbResources 個資源須屬 同一 GPU device。
• 多個 SM 型資源組合時，須來自 同一次 split 呼叫 且有 相同 coscheduledSmCount（若非 remainder 部分）。
• 最多只能有一個 workqueue config 或 workqueue 型資源。
• 被組合的資源須在 resources 陣列中 連續配置。可同時組合不同型（如 SM + WQ）。

Step 4：建立 Green Context

cudaError_t cudaGreenCtxCreate(cudaExecutionContext_t* phCtx,
                               cudaDevResourceDesc_t desc,
                               int device, unsigned int flags);

cudaSetDevice(current_device);            // 或 cudaInitDevice，先初始化 primary context
cudaExecutionContext_t green_ctx{};
cudaGreenCtxCreate(&green_ctx, resource_desc, current_device, 0);

• flags 應設 0。GC 只能存取 descriptor 內封裝的資源，這些資源於此步被 provisioned。
• 建議 先以 cudaInitDevice 或 cudaSetDevice 顯式初始化該 device 的 primary context，避免建立 GC 時額外的 primary context 初始化開銷。
• 建立成功後，可對該 EC 逐型呼叫 cudaExecutionCtxGetDevResource 驗證其資源。

多 GC 與 oversubscription：多數情況各 GC 有 非重疊 的 SM 集合（如 GC1 用 result[2..4]、GC2 用 result[0..1]，每顆 SM 只屬一個 GC）。但 SM oversubscription 也可能（如 GC2 用 result[0..2]，則 result[2] 的 SM 被兩個 GC 共用）。

Launching Work

在 GC 上啟動 kernel，先用 cudaExecutionCtxStreamCreate 為該 GC 建立 stream，於其上 <<<>>> 或 cudaLaunchKernel 啟動的 kernel 只能用該 stream 的 EC 可用資源：

cudaStream_t green_ctx_stream;
int priority = 0;
cudaExecutionCtxStreamCreate(&green_ctx_stream, green_ctx,
                             cudaStreamDefault, priority);
my_kernel<<<grid_dim, block_dim, 0, green_ctx_stream>>>();
cudaGetLastError();

• 當 green_ctx 是 green context 時，傳入的預設 stream 建立 flag 等同 cudaStreamNonBlocking。
• 只要使用者把屬於 GC 的 stream 傳給 library，library 也能輕鬆受惠於 GC。

CUDA Graphs 的細節

• stream capture：參與 capture 的 stream 之 EC 決定相關 graph node 的 EC。
• graph API：使用者須為每個相關 node 顯式設定 EC。加 kernel node 應用多型的 cudaGraphAddNode（type cudaGraphNodeTypeKernel），並設 cudaKernelNodeParamsV2 之 .kernel.ctx 欄位；應避免 cudaGraphAddKernelNode（不允許指定 EC）。
• 同一 graph 內不同 node 可屬不同 EC。
• 驗證：Nsight Systems node tracing mode（--cuda-graph-trace node）可觀察各 node 實際執行的 GC（預設 graph tracing mode 會把整個 graph 歸在 launch stream 的 GC 下，無法反映各 node 的 EC）。程式化驗證可用 driver API cuGraphKernelNodeGetParams 比對 node_params.ctx（CUgraphNode 與 cudaGraphNode_t 可互換，但需 include cuda.h 並 -lcuda）。

Thread Block Clusters

帶 thread block cluster 的 kernel 可像一般 kernel 一樣在 GC stream 上啟動，用該 GC 的 provisioned 資源。Step 2 已說明 split 時如何指定 coscheduled SM 數以利 cluster。仍應用 occupancy API 求最大可行 cluster 大小：

cudaLaunchConfig_t config = {0};
config.gridDim = grid_dim;          // 須為 cluster dim 的倍數
config.blockDim = block_dim;
config.dynamicSmemBytes = expected_dynamic_shared_mem;
cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeClusterDimension;
attribute[0].val.clusterDim = {1, 1, 1};
config.attrs = attribute; config.numAttrs = 1;
config.stream = gc_stream;          // 必須傳該 kernel 會用的 stream
int max_potential_cluster_size = 0;
cudaOccupancyMaxPotentialClusterSize(&max_potential_cluster_size, cluster_kernel, &config);
int num_clusters = 0;
cudaOccupancyMaxActiveClusters(&num_clusters, cluster_kernel, &config);

• 把 GC stream 設為 cudaLaunchConfig 的 stream，則 cudaOccupancyMaxPotentialClusterSize 與 cudaOccupancyMaxActiveClusters 會 把該 GC provisioned 的 SM 納入考量。
• 此用例對「被傳入 GC stream 的 library」與「GC 由 remaining 資源建立」的情況特別重要。

驗證 GC 使用

• Nsight Systems：不同 GC 的 kernel 出現在 CUDA HW timeline 的不同 Green Context row。
• Nsight Compute：Session 頁的 Green Context Resources 提供 provisioned 資源的 bitmask（可確認 GC 間無重疊或預期的 oversubscription 重疊）；Details 的 Launch Statistics 列出該 GC provisioned 的 SM 數。

Additional Execution Context APIs

完整範例（critical kernel 提早完成）

情境：兩個 kernel 跑在兩條 non-blocking stream。先在 strm1 啟動長跑 kernel delay_kernel_us（block 數超過 SM 數、需多波），短暫等待後在 strm2 啟動短而關鍵的 critical_kernel。量測兩者 GPU duration 與「CPU 啟動到完成」時間。

無 GC（critical 在高 priority stream，但仍須等部分 block）
stream13 ████████ delay_kernel ███████████████████
stream14         ⟵ 等 0.9ms ⟶ ███ critical ███          ← 被延遲

有 GC（GC2=112 SM 給 delay，GC3=16 SM 給 critical，無重疊）
GC2(112SM) ████████ delay_kernel ██████████████████████ ← 變慢
GC3( 16SM) ███ critical ███                                ← 幾乎即時開始

• 不用 GC、只靠高 priority stream：critical kernel 可在 delay kernel 部分 block 完成後開始，但仍須等某些長 block，故被延遲（原文範例為 0.9ms；若兩 stream 同 priority 會更晚）。
• 用兩個非重疊 GC（H100 132 SM，範例切 16 SM 給 critical、112 SM 給 delay）後，critical kernel 幾乎即時開始。
• 代價：critical 與 delay kernel 各自 duration 可能 變長（受限於各 GC 的 SM）。但 關鍵結果 是 critical work 能 顯著提早開始與完成。實際 SM split 應逐案實驗決定。

考試/測驗重點

Related Notes

• 04-CUDA-Features/03-CUDA-Graphs-Structure-and-Capture
• 04-CUDA-Features/04-CUDA-Graphs-Updating-and-Conditional
• 04-CUDA-Features/07-Cooperative-Groups-Deep-Dive
• 04-CUDA-Features/08-Programmatic-Dependent-Launch-Deep-Dive
• 04-CUDA-Features/16-Work-Stealing-Cluster-Launch-Control
• 04-CUDA-Features/25-Driver-Entry-Point-Access
• 03-Advanced-CUDA/01-Advanced-Launch-and-Clusters
• 03-Advanced-CUDA/08-CUDA-Driver-API
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 02-Programming-GPUs/15-Async-Callbacks-Ordering-Graphs
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps