CUDA Dynamic Parallelism (動態平行) 練習題 (Practice - Dynamic Parallelism Fundamentals and Overview)

CDP 是 CUDA 程式模型的擴充，讓 一個 kernel 內部可以 launch 另一個 kernel，使在 device 上執行的 thread 能自行產生新的 grid。早期 CUDA 中 grid 只能由 host code launch。最受益的演算法特徵：recursion、irregular loop structures、time-space variation（時空變化）、adaptive grid refinement——這些不適合「平坦且單層」平行度的工作，過去得靠多次 host kernel call 或迴圈序列化，增加 host 負擔與 host-device 通訊。

語法與 host 端完全相同：kernel<<<Dg, Db, Ns, S>>>(args)。

• Dg (dim3)：grid 維度與大小（block 數）。
• Db (dim3)：每個 block 的維度與大小（thread 數）。
• Ns (size_t)：每 block 動態配置的 shared memory bytes，額外於靜態配置，預設 0。
• S (cudaStream_t)：此呼叫關聯的 stream，預設 0 (NULL stream)。

S 限制：stream 必須在發起呼叫的同一個 thread block 內建立，其 scope 為 block-private。

兩個問題：(1) 錯失平行度——若迴圈體本可平行，卻被寫成序列執行；(2) control divergence——同一 warp 內各 thread 的迴圈次數差異大時，warp 必須等最長的 thread，閒置浪費。
CDP 解法（Fig. 21.5，拆成 parent + child）：迴圈迭代改由 child kernel threads 並行執行（抽取更多平行度）；每個 child thread 只做一次迭代 → load balance 佳、消除 divergence。

(1) 索引單位：parent 用 一條 thread 處理一組控制點（索引以 thread 為基底 threadIdx.x+blockDim.x*blockIdx.x），原版用一個 block（索引以 blockIdx.x）；因為 parent 工作（算 curvature + launch）很小。
(2) device-side cudaMalloc：vertexPos 改成指標，在 kernel 內依實際 nVertices 精準配置記憶體，避免原版每條線固定配 MAX_TESS_POINTS（最壞情況）造成的浪費。
(3) 釋放：device 端配的記憶體必須由 device kernel cudaFree 釋放，故另寫 freeVertexMem kernel 由 host 呼叫平行釋放（host 不能直接 free device 端 cudaMalloc 的記憶體）。

可傳：global、constant、texture memory（皆 device 全域共享）。不可傳：local memory（thread 私有）、shared memory（block 私有）——傳了會讓 child 存取無效位址。
兩個一致視圖時刻：(1) parent launch child 那一刻——launch 之前的所有寫入 child 一定看得到，launch 之後的寫入無保證；(2) parent 同步於 child 完成那一刻——在此之前 child 的寫入不保證對 parent 可見。心法：要給 child 讀的就在 launch 前寫完；要讀 child 結果就先同步再讀。

Pending launch pool 是追蹤「執行中或等待執行」kernel 的固定大小緩衝區，預設可容納 2048 個 pending kernel calls。超過 → 啟用 virtualized pool，效能掉一個數量級（10× 以上）。
調整：在 host 端、launch parent 之前呼叫 cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, N)，把固定 pool 設為「預期 launch 的 grid 數」。一般建議：只在預期 launch 數超過預設 2048 時才調高；調得比實際需求小（如 1024）既無加速也無意義。

不指定 stream → 使用 default NULL stream，同一 block 內所有 launch（即使來自不同 thread）都被序列化。stream 的 scope 是 block-private（不可跨 block 使用別 block 建的 stream）。
要併發須讓每個 thread 建立自己的 non-blocking named stream：

cudaStream_t s;
cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
child<<<gridDim, blockDim, 0, s>>>(...);
cudaStreamDestroy(s);

否則 Bezier 這類「同 block 多 thread 各 launch child」的程式平行度甚至會比沒用 CDP 還差。

Nesting depth = child kernel 再 launch kernel、層層相疊所達到的總層數，目前硬體上限為 24 層。像 quadtree 這種遞迴 kernel launch 前必須先檢查（if (depth >= max_depth) return;），否則超過上限會 launch 失敗。
Synchronization depth：當存在 parent-child 同步時，系統需保存 parent grid 的狀態，記憶體成本更高，因此對巢狀深度施加額外、更嚴格的限制（≤ nesting depth）。

(a) False。__constant__ 是編譯期 file-scope 宣告，不是 runtime 動態建立的，parent 無法在執行期「定義新的」constant 變數給 child 繼承。
(b) False。shared memory 是 block 私有、local memory 是 thread 私有；parent 不可把它們的指標傳給 child grid，child 也無法存取（會是無效位址）。child 只能存取 global / constant / texture memory。

(a) False。每條 parent thread 各 launch 一個 child grid → child kernel 數 = N_LINES = 1024，不是 16；16 只是 parent block 數 = N_LINES/BLOCK_DIM = 1024/64。
(b) False。預設 2048 已足夠涵蓋 1024 個 launch；調小不會變快，反而可能限制。只有「預期 launch 數 > 2048」才該調高。
(c) False。per-thread stream 是每條 thread 各建一個 → 共 1024 個 stream，不是 16。

點數每深一層除以 4：64 → 16 → 4 → 1。
(a) 最大深度 = 4（depth 0, 1, 2, 3）。depth 0 有 64 點、depth 1 每節點 16 點、depth 2 每節點 4 點（>2 仍切）、depth 3 每節點 1 點（≤2 停）。
(b) depth 0、1、2 的所有 active 節點都會 launch（每次 launch 一個 <<<4>>> child grid），節點數為 1、4、16 → 總 child kernel launch = 1 + 4 + 16 = 21 = (4³−1)/3。（depth 3 各節點只剩 1 點，不再 launch。）

6 個。stream scope 是 block-private，NULL stream 使同一 block 內所有 launch 序列化，所以每個 block 同時最多只有 1 個 child grid 在跑；6 個 block 彼此獨立可併發 → 最多 6 × 1 = 6 個 child kernel 併發。（若改用 per-thread named streams，則同 block 內的 child grid 才能併發，上限大幅提高。）

for-loop 版：迴圈被序列執行 → 錯失平行度；warp 內迭代數不一 → control divergence；block 間迭代數差異大 → 部分 SM 提早閒置、利用率下降。
child-grid 版：迴圈迭代由 child threads 並行（抽取更多平行度）；每 child thread 只做 1 次迭代 → load balance 佳、消除 divergence、SM 利用率提升。
CDP 版反而更慢的情況：(1) child grid 太小（thread 太少）→ 嚴重 GPU underutilization，launch overhead 蓋過收益；(2) 沿用 default NULL stream → 同 block 的 child grid 全部序列化，平行度甚至比原版更差，須改用 per-thread named streams；(3) launch 數超過 pending pool（2048）落入 virtualized pool，慢一個數量級。結論：CDP 收益需「夠胖的 child grid + named streams + 足夠的 pool」才能兌現。

為何 scatter：reorder 的目標是把任意分布的輸入點依所屬象限動態分群到輸出 buffer，每個象限的輸出長度執行期才知道、且要連續排列。gather 要求每個 thread「固定知道自己要讀哪個來源」，但這裡是輸入決定輸出位置（一對多分群），無法事先靜態映射 → 必須 scatter：對 smem2[q] 做 atomicAdd 動態取得該象限「下一個可寫位置」。
代價：需要 atomic（搶象限內下一格），有競爭成本；寫出位置不規則，較難 coalesce。但因計數器只有 4 個且在 shared memory，atomic 競爭遠比 global 便宜。
同構：count（shared-mem atomic）→ 4-element scan（求 offset）→ scatter（reorder） 與 radix sort 的 count → scan → scatter 三步完全同構。

元素只有 4 個，序列 scan 僅需 3 次加法、無 barrier、無 SIMD 浪費；若改用平行 Kogge-Stone，反而要付出 __syncthreads()、額外 shared memory、warp 內多數 lane 閒置等 overhead，對這麼小的問題完全不划算。平行 scan 的 work（O(N log N)）與同步成本只有在 N 很大時才回本。
這呼應「演算法選擇要看資料規模」的原則：最佳演算法依輸入大小而異——小資料用最簡單的序列法、大資料才用平行法。與 Kogge-Stone vs Brent-Kung 的取捨（work-efficiency vs 步數）同屬「沒有放諸四海皆準的最佳解」這一類。

核心心法：Dynamic Parallelism 把「動態發現工作 → 立即就地平行化」的迴路留在 device 上，藉此抽取更多平行度、消除 warp 內 control divergence、並用 device-side cudaMalloc 精準配置記憶體。但它不是免費午餐——收益必須靠「夠胖的 child grid + per-thread named streams + 足夠的 pending pool + 受控的 nesting depth」才能兌現；否則小 grid 的 underutilization、NULL stream 的序列化、virtualized pool 的退化，都可能讓 CDP 版比原始版本更慢。