重要執行考量與總結

重點總覽 (Overview)

要正確且有效地使用 Dynamic Parallelism,必須掌握四大執行考量。下表為總覽:

記憶體與資料可見性 (Memory and Data Visibility)

當 parent thread 把指標傳給 child grid 時,必須保證該記憶體對 child 是可存取的,且要清楚何時寫入的資料才會彼此可見。

可傳遞 (兩者皆可存取):

兩個一致視圖時間點 (Two coherence points)

parent 與 child 對記憶體只有兩個保證一致的時刻:

parent thread 時間軸 (memory writes)
  │
  ▼  write A  ──────────► [可見] child 看得到 launch 前的所有寫入 (A)
  │
 ┌┴───────────────┐  ← (1) LAUNCH child grid   ← 一致視圖點 #1
 │  launch child   │
 └┬───────────────┘
  │  write B  ──────────► [不保證] child 不一定看得到 launch 後的寫入 (B)
  │
  │     ┌──────────── child grid 執行中, write C ────────────┐
  │     └──────────────────────────────────────────────────┘
  │
 ┌┴───────────────┐  ← (2) parent SYNC 於 child 完成 ← 一致視圖點 #2
 │  cudaDevice...  │
 └┬───────────────┘
  ▼  read C  ──────────► [可見] 同步後才保證看得到 child 的寫入 (C)

• 點 #1 — child launch 時刻:launch 之前的所有寫入,child 一定看得到;launch 之後的寫入,無保證。
• 點 #2 — parent 同步於 child 完成時:在此之前,child 的寫入不保證對 parent 可見。

Pending Launch Pool 配置 (Pending Launch Pool Configuration)

Pending launch pool = 追蹤「執行中或等待執行」kernel 的緩衝區,大小固定。

                ┌─────────────────────────────┐
launches ─────► │  Fixed pool (預設 2048 slots) │ ──► 快速排程
                └─────────────────────────────┘
                          │ 超過 2048
                          ▼
                ┌─────────────────────────────┐
                │   Virtualized pool          │ ──► 慢 10× 以上
                └─────────────────────────────┘

• 預設可容納 2048 個 pending kernel calls。
• 超過 → 啟用 virtualized pool,效能掉 一個數量級 (order of magnitude) 以上。
• 解法:在 host 端、launch parent 之前呼叫:

// 把固定 pool 調到「預期 launch 的 grid 數」
cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, N_LINES);
computeBezierLines_parent<<<N_LINES/BLOCK_DIM, BLOCK_DIM>>>(...);

Bezier 範例量化:

Per-Thread Streams (Streams)

device thread 也能像 host 一樣用 streams 讓 child grids 併發。關鍵:stream scope 是 block-private。

同一個 block (多個 thread 各 launch 一個 child grid)

【NULL stream — 序列化】          【Named streams — 併發】
 t0 ▓▓▓                            t0 ▓▓▓
 t1     ▓▓▓                        t1 ▓▓▓
 t2         ▓▓▓                    t2 ▓▓▓
 t3             ▓▓▓                t3 ▓▓▓
 ───────────時間──────►            ───時間──►
 (浪費 SM,平行度大降)             (充分利用 SM)

取代 Bezier 範例 (Fig. 21.7) line 19 的程式碼 (Fig. 21.12):

cudaStream_t stream;
// 每個 thread 建立自己的 non-blocking stream (block-private)
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
computeBezierLines_child<<<gridDim, blockDim, 0, stream>>>(
        bLines_d, lidx, bLines_d[lidx].nPoints);
cudaStreamDestroy(stream);

巢狀深度與同步深度 (Nesting Depth & Synchronization Depth)

child kernel 自己也能再 launch kernel,層層相疊。

depth 0  parent grid (host launch)
   └─► depth 1   child grid
          └─► depth 2   grandchild grid
                 └─►  ...           ...
                        └─► depth 23  (最深合法層)
                               ✗ depth 24+ → 不合法

• Nesting depth:總共抵達的層數;目前硬體上限 = 24 層。
• 像 quadtree 這種遞迴,launch 前必須先檢查是否已達 max depth (if (depth >= MAX_DEPTH) return;),否則 launch 失敗。
• Synchronization depth:當存在 parent-child 同步時,系統需保存 parent grid 狀態,記憶體成本更高 → 對巢狀深度有額外更嚴格的限制 (≤ nesting depth)。

本章總結 (Chapter Summary)

Dynamic Parallelism 擴充 CUDA 模型,讓 kernel 可呼叫 kernel,核心價值:

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

Related Notes

• 21-CUDA-Dynamic-Parallelism/01-Dynamic-Parallelism-Fundamentals
• 21-CUDA-Dynamic-Parallelism/02-Bezier-Curves-Example
• 21-CUDA-Dynamic-Parallelism/03-Recursive-Quadtree-Example
• 20-Heterogeneous-Computing-Cluster/03-Overlapping-Computation-and-Communication
• 04-Compute-Architecture-And-Scheduling/03-Resource-Partitioning-and-Occupancy
• 05-Memory-Architecture-And-Data-Locality/01-Memory-Access-Efficiency-and-CUDA-Memory-Types