Kernel 執行控制 (Kernel Execution Control)

重點總覽 (Overview)

本節整理 CUDA 在「kernel 執行期能力」上的演進，重點不在效能，而在軟體工程能力 (composability、可維護性、除錯) 與排程彈性 (priority、load balance)。

Kernel 內函式呼叫 (Function Calls within Kernel Functions)

早期模型：原始碼可以「看起來」有函式呼叫，但 compiler 必須把所有函式 body inline 進 kernel object，runtime 時 kernel 內沒有真正的 call。

• 不支援：system calls、dynamically linked library calls、recursion、virtual functions (C++)。
• 對 performance-critical 區段尚可，但無法支撐複雜應用的軟體工程實務。

Kepler / CUDA 5 之後：支援 runtime 真正的函式呼叫。

• Compiler 不再被強制 inline (仍可為了效能選擇 inline)。
• 由 per-thread 的 massively parallel call frame stack (有 cache、快速實作) 支撐。
• 帶來 composability：不同作者寫不同 kernel 元件再組裝，免大改；廠商可釋出不含原始碼的 device library (IP 保護)。

// CUDA 5+ (Kepler): 真正的 runtime function call，不再強制 inline
__device__ int factorial(int n) {      // recursion 現在合法
    if (n <= 1) return 1;
    return n * factorial(n - 1);       // 遞迴呼叫使用 per-thread call stack
}

__global__ void demo(int *out, int n) {
    out[threadIdx.x] = factorial(n);
    printf("tid %d done\n", threadIdx.x);  // 在 kernel 內呼叫標準函式庫
}

• 支援 recursion，並能直接「cut-and-paste」CPU code 取得堪用 kernel (後續仍建議調效)。
• kernel 可呼叫 printf()、malloc()。
  • printf() 對 production 除錯特別有用：可加一個 dump 內部狀態的模式，讓非技術 end user 也能回報有意義的 bug report。

Device lambdas (CUDA 8 + C++11)

• 搭配 metaprogramming → 高效能、可重用 code。
• lambda 可作為參數傳入 kernel：寫泛型 kernel，例如 sorting kernel 把比較函式當輸入參數。
• Extended lambdas (--extended-lambda 旗標)：可在 C++ lambda 上加 __host__ __device__，進一步簡化可重用 code。

// 由比較 lambda 參數化的泛型 kernel
template <typename Compare>
__global__ void sort_kernel(int *data, int n, Compare cmp) {
    // ... 用 cmp(a, b) 決定順序 ...
}

// extended lambda：同一份程式碼可在 host / device 執行
auto cmp = [] __host__ __device__ (int a, int b) { return a < b; };
sort_kernel<<<grid, block>>>(d_data, n, cmp);

Kernel 內例外處理 (Exception Handling in Kernel Functions)

• 早期 CUDA：kernel code 完全不支援 exception handling。
• 對 performance-critical 區段影響不大，但 production 品質應用常依賴 exception 來偵測/處理罕見情況 (不必額外寫顯式測試碼)。
• 目前提供有限 (limited) 的 exception handling，主要透過 CUDA debugger 體現：

多 Grid 同時執行 (Simultaneous Execution of Multiple Grids)

• 最早期：每個 device 同一時間只能執行一個 grid。多個 grid 雖可用 CUDA streams 提交，但被排在 queue 中緩衝，前一個跑完才放下一個。
• Fermi 之後：允許同一應用的多個 grid 同時執行。
  • 降低開發者把多個 kernel「batch」成一個大 kernel 以填滿 device 的壓力。
  • 可把工作切成不同 priority 的 chunk。

Pre-Fermi (一次一個 grid，串列):
  stream queue : [Grid A][Grid B][Grid C]
  GPU 時間軸   : |==A==|==B==|==C==|     小 grid 無法填滿 SM -> 低利用率

Fermi+ (多 grid 同時執行):
  SMs : |===== Grid A (low-prio local) =========|
        |   |== Grid B (high-prio remote) ==|   |   <- 立刻插入執行
        time -->

典型情境 (parallel cluster)：工作分成 local 與 remote partition，remote 牽涉與其他 node 互動、位於全域進度的 critical path 上。

硬體佇列與動態平行 (Hardware Queues and Dynamic Parallelism)

Kepler / CUDA 5：在多 grid launch 之上加入多條 hardware queues，更有效率地排程「來自多個 stream 的多個 grid」的 thread blocks。

單一 HW queue (false serialization):
  stream1: A1->A2     stream2: B1
  HW queue: [A1][A2][B1]   <- B1 雖獨立卻被卡在 A2 之後

多 HW queues (Kepler):
  stream1 --> HWQ0: [A1][A2]
  stream2 --> HWQ1: [B1]      <- B1 與 A1/A2 併發派發

Dynamic parallelism (詳見 21-CUDA-Dynamic-Parallelism/01-Dynamic-Parallelism-Fundamentals)：

• GPU grid 可非同步、動態、依資料/計算負載地 launch child grids (GPU work creation)。
• 減少 CPU-GPU 互動與同步；CPU 因此可去做其他有用運算。

可中斷的 Grid (Interruptable Grids)

• Fermi：允許正在執行的 grid 被「cancel」。
• 讓 CUDA 加速應用可隨時中止長時間計算，且開發者幾乎不需特別設計。
• 用途：
  • 實作 user-level task scheduling 系統。
  • GPU node 間更好的 load balance。
  • 某顆 GPU 負載過重、跑得比同儕慢時，能更優雅地處理。

協作式 Kernel (Cooperative Kernels)

GPU 處理 irregular data 常受 load imbalance 之苦。CUDA 11 引入 cooperative kernels。

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void coop_kernel(/* ... */) {
    cg::grid_group grid = cg::this_grid();
    // ... phase 1 ...
    grid.sync();   // grid-wide barrier：唯有「所有 block 併發」才安全
    // ... phase 2 ...
}

// 啟動：grid 必須能塞進 GPU，不可 oversubscribe
void *args[] = { /* ... */ };
cudaLaunchCooperativeKernel((void*)coop_kernel, gridDim, blockDim, args);

一般 grid (block 數 > SM 容量) -> 分波 (waves)，排程順序未定義
  若呼叫 grid-wide barrier -> 等待中的 block 永遠不會上場 -> DEADLOCK

Cooperative grid (block 數 <= GPU 容量) -> 全部同時駐留
  SM0:[B0][B1]   SM1:[B2][B3]  ... 全在跑  -> grid.sync() 安全

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 22-Advanced-Practices-And-Future-Evolution/01-Host-Device-Interaction-Memory-Model
• 22-Advanced-Practices-And-Future-Evolution/03-Memory-Bandwidth-and-Compute-Throughput
• 21-CUDA-Dynamic-Parallelism/01-Dynamic-Parallelism-Fundamentals
• 21-CUDA-Dynamic-Parallelism/04-Execution-Considerations-and-Summary
• 20-Heterogeneous-Computing-Cluster/03-Overlapping-Computation-and-Communication
• 04-Compute-Architecture-And-Scheduling/01-GPU-Architecture-and-Block-Scheduling