Enter your search text in the box above

Select a result to preview

Dynamic Parallelism 基礎與概觀

#advanced-practice #dynamic-parallelism #parent-child-kernels #load-balancing #control-divergence

重點總覽 (Overview)

CUDA Dynamic Parallelism (CDP) 讓 kernel 內部可以再 launch 另一個 kernel,使裝置上執行的 thread 能夠自行產生新的 grid。在早期 CUDA 中,grid 只能由 host code 啟動。

項目 沒有 Dynamic Parallelism 有 Dynamic Parallelism
誰能 launch grid 只有 host (CPU) host 與 device thread 皆可
動態發現的新工作 grid 終止 → 回報 host → host 再 launch thread 直接就地 launch child grid
host-device 通訊負擔 高 (每次都要往返) 低
不規則/遞迴/時空變化工作 需多次 host kernel call 或迴圈序列化 自然以 parent/child 表達
典型受益場景 — recursion、irregular loops、time-space variation、adaptive refinement
Important

CDP 的核心價值不只是「方便」:它把動態發現工作 → 立即平行化的迴路留在 device 上,藉此 擷取更多平行度 並 消除 warp 內的 control divergence。

為何需要 Device-launched Grids (Background)

許多真實應用的工作量會 隨空間變化 或 隨時間動態變化。書中以 turbulence simulation (亂流模擬) 為例:燃燒流由左往右,右側活動劇烈、需要更細的網格;左側平靜、粗網格即可。

網格策略 左側 (低細節區) 右側 (高細節區) 缺點
Fixed fine grid 過量運算、無增益 精度足夠 浪費資源 / 多餘資料
Fixed coarse grid 剛好 精度不足 犧牲準確度
Dynamic grid 維持粗網格 偵測到劇變後就地 refine 需 device 端能動態 launch 
空間變化的工作量 (turbulence)
  低細節 ───────────────► 高細節
  ┌───────────────┬──────┬──┬─┐
  │               │      │░░│░│   ░ = 需要 refine 的區域
  │   coarse      │      ├──┼─┤
  │               │      │░░│░│
  └───────────────┴──────┴──┴─┘
  理想: 只在劇變區投入更多運算資源 (fine grid)

Grid launch 模式對比 (Fig. 21.2) — 時間軸由上往下:

(A) 無 Dynamic Parallelism            (B) 有 Dynamic Parallelism
   HOST            GPU                    HOST            GPU
    │  launch ─►  grid0                    │  launch ─►  grid0
    │            (發現新工作)               │            (thread 發現新工作)
    │ ◄─ 終止+回報                          │              │ launch ─► child0
    │  launch ─►  grid1                    │              │ launch ─► child1
    │ ◄─ 終止+回報                          │              ▼ (不需回到 host)
    │  launch ─►  grid2                    │            完成
   每次新工作都要一次 host↔device 往返       device 端就地展開,host 不介入
Tip

(A) 中後續 grid 雖畫成依序啟動,但 host 仍可用不同 stream 或合併讓它們並行 —— 重點是仍需 host 介入;(B) 完全免去這個往返。

Parent / Child Kernel 模型 (Dynamic Parallelism Overview)

從程式設計者角度:可以在一個 kernel 內寫一行呼叫另一個 kernel function。下例中 host 啟動 A、B、C 三個 kernel(一如全書慣例);新東西是 kernel B 又啟動了 X、Y、Z(在不支援 CDP 的早期 CUDA 這是非法的)。

   CPU (host)            GPU (device)
   main()
     ├── A <<<>>>  ─────►  A
     ├── B <<<>>>  ─────►  B ──┬── X <<<>>>  ─► X   ← kernel 內 launch kernel
     │                         ├── Y <<<>>>  ─► Y
     │                         └── Z <<<>>>  ─► Z
     └── C <<<>>>  ─────►  C

Kernel-in-kernel Launch 語法

device 端呼叫 kernel 的語法 與 host 端完全相同:

kernel<<< Dg, Db, Ns, S >>>(args...);
參數 型別 意義 預設
Dg dim3 grid 的維度與大小 (block 數) 必填
Db dim3 每個 thread block 的維度與大小 (thread 數) 必填
Ns size_t 每個 block 動態配置 的 shared memory bytes(額外於靜態配置) 0
S cudaStream_t 此次呼叫關聯的 stream 0 (NULL stream)
Warning

S 指定的 stream 必須在發起呼叫的同一個 thread block 內建立;stream 的作用域 (scope) 私有於該 block。若省略 S,block 內所有 thread 共用 default NULL stream,導致 同一 block launch 的 child grid 全部序列化(細節見 21-CUDA-Dynamic-Parallelism/04-Execution-Considerations-and-Summary)。

從迴圈到 Child Grid:消除 Divergence (The Loop-to-Child Transformation)

書中用一個「不計算有用結果、但結構常見」的假想 kernel 來對比。

Fig. 21.4 — 沒有 Dynamic Parallelism:每個 thread 先做一段運算,再 迴圈 處理自己負責的一串資料元素。

__global__ void kernel(unsigned int* start, unsigned int* end,
                       float* someData, float* moreData) {
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    doSomeWork(someData[i]);                 // 每個 thread 的前置運算
    for (unsigned int j = start[i]; j < end[i]; ++j) {
        doMoreWork(moreData[j]);             // 逐元素、序列、長度因 thread 而異
    }
}

此結構在多處反覆出現:

應用 前置工作 (doSomeWork) 迴圈工作 (doMoreWork)
Graph search 拜訪一個 vertex 逐一走訪鄰居 vertex(類似 vertex-centric BFS)
Sparse matrix 找出某 row 的起點 逐一處理該 row 的 nonzero
Adaptive simulation 處理一個 coarse grid 元素 逐一處理需細分的 finer grid 元素

兩個問題:

  1. 錯失平行度:若迴圈 (line 07-09) 的工作本可平行,卻被寫成序列。
  2. Control divergence:若同一 warp 內各 thread 的迴圈次數差異大,divergence 會拖垮效能。
warp 內 4 個 thread 的迴圈長度差異 (X = 一次 doMoreWork)
  t0: X X
  t1: X X X X X X X X      ← warp 必須等最長的 t1
  t2: X
  t3: X X X
       └── t0/t2/t3 在尾段閒置 (divergence)

Fig. 21.5 — 有 Dynamic Parallelism:原 kernel 拆成 parent + child。parent 不再迴圈,而是 launch 一個 child grid 把迴圈體的工作平行展開。

__global__ void parentKernel(unsigned int* start, unsigned int* end,
                             float* someData, float* moreData) {
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    doSomeWork(someData[i]);                 // 與原 kernel 相同的前置
    // 以「每 thread 負責的元素數」決定 child grid 大小
    childKernel<<< ceil((end[i]-start[i])/256.0), 256 >>>
                  (start[i], end[i], moreData);
}

__global__ void childKernel(unsigned int start, unsigned int end,
                            float* moreData) {
    unsigned int j = start + blockIdx.x*blockDim.x + threadIdx.x;
    if (j < end) {
        doMoreWork(moreData[j]);             // 原本迴圈體 → 現在一個 thread 做一次
    }
}

效益對應問題:

原問題 CDP 的解法
錯失平行度 迴圈迭代改由 child threads 並行 執行
Control divergence 每個 child thread 只做 一次 迭代 → load balance 佳、無 divergence
Tip

同樣效果也能用人工改寫(如 edge-centric BFS)達成,但對某些應用而言這類轉換 笨拙、複雜且易錯。CDP 提供一個自然的表達方式。

Child grid 大小公式(每個 parent thread 各自決定):

gridDim  = ceil( (end[i] - start[i]) / BLOCK_DIM )   // 本例 BLOCK_DIM = 256
blockDim = BLOCK_DIM
Warning

CDP 不是「越多 child grid 越好」。若 child grid 只有極少 thread,會嚴重 underutilize GPU。一般建議 child grid 要有 夠多 block,或至少 block 內有夠多 thread(見總結 21-CUDA-Dynamic-Parallelism/04-Execution-Considerations-and-Summary)。

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

情境 / 關鍵字 答案 / 技巧
「early CUDA 中 grid 只能由誰 launch?」 只能由 host code launch;CDP 才允許 device thread launch
「kernel 內呼叫 kernel 的語法?」 與 host 相同:kernel<<<Dg, Db, Ns, S>>>(...)
「Ns / S 各是什麼?」 Ns=動態 shared memory bytes (預設 0);S=stream (預設 NULL,須在同 block 建立)
「無 CDP 時新工作如何處理?」 grid 終止 → 回報 host → host 重新 launch(一次往返)
「parent grid vs child grid 誰做什麼?」 parent 做前置並 launch;child grid 平行執行原迴圈體
「Fig. 21.4 的兩個問題?」 (1) 錯失平行度 (2) warp 內迴圈長度不一 → control divergence
「CDP 如何消除 divergence?」 每 child thread 只做一次迭代 → load balance、無 divergence
「適合 CDP 的演算法特徵?」 recursion、irregular loops、time-space variation、adaptive grid refinement
「default NULL stream 的陷阱?」 同 block 內所有 launch 序列化;要並行須用 named per-thread streams
「CDP 的反效果?」 child grid thread 太少 → GPU underutilization