第四章練習題 (Practice - CUDA Features)

作答前可先用這張表喚起記憶；每列是「情境/關鍵字 → 答案」。

hardware-coherent = CPU/GPU 共用「邏輯合併的 page table」（如 Grace Hopper），一致性粒度為 cache-line；software-coherent = CPU/GPU 各自獨立 page table，以 page fault + migration 模擬一致性，粒度為整頁。
兩大優勢：更少 page fault（不需以 fault 模擬一致性或搬移記憶體）、更少 contention（粒度為 cache-line 而非整頁，同頁不同 cache-line 不互相干擾）。

因為任何 kernel 執行期間 GPU 獨佔「所有」managed data（即使該 kernel 沒用到那塊），CPU 同時存取會被視為對 CPU 不可存取而 segfault。
正解是先 cudaDeviceSynchronize（或任何邏輯上保證 GPU 完工的呼叫）再存取；更細粒度替代是用 cudaStreamAttachMemAsync 把配置綁到特定 stream，把獨佔從 whole-GPU 縮為 per-stream。

三階段：definition（建 cudaGraph_t 的節點與依賴）→ instantiation（cudaGraphInstantiate 取快照、驗證、產出 executable cudaGraphExec_t）→ execution（cudaGraphLaunch 送入 stream，可重複多次而不需重新 instantiate）。
capture 不可用於 cudaStreamLegacy（NULL stream），但可用 cudaStreamPerThread。EndCapture 必須在 origin stream（呼叫 BeginCapture 那條）呼叫，其他因 event 依賴加入的 stream 都須 join 回 origin，否則 capture 失敗。

拓撲與節點類型不變、只改節點參數（kernel 參數、記憶體位址）時可用 graph update，更新於「下一次 launch」生效、不影響進行中的 launch；拓撲或節點類型改變則必須 re-instantiate。
IF：condition 非零執行 body 一次（size==2 才有 else，於 condition 為零時執行）；WHILE：非零就反覆執行 body，進入時評估一次、每次 body 完成後再評估；SWITCH：執行第 n 個 zero-indexed body，condition 不對應任何 body 則不執行（非 fallthrough）。

因為 graph allocation 的「虛擬位址」在整個 graph 生命固定，圖內其他操作可直接引用該 VA；底層 physical memory 變動由 driver 重新映射、與 VA 解耦，故不需 graph update。
圖內：把同一 VA range 指派給生命期「不重疊」的不同配置（故不相交配置的指標不保證唯一）；圖間：用 virtual aliasing 把同一 physical memory 映射到各 graph 獨有的 VA（多個不會同時執行的 graph 共用同一份 physical）。

cudaMalloc / cudaFree 會讓 GPU 跨「所有」stream 同步；stream-ordered allocator 把配置/釋放排入 stream，與 kernel、async copy 一起依序執行，不阻塞 host 或其他 stream。
跨 stream 時須用 event 或 stream 同步（cudaEventRecord + cudaStreamWaitEvent）保證「存取在配置之後、free 在所有存取完成之後」，否則為未定義行為（free 開始後再使用 allocation 同樣未定義）。

四個 accessor：this_thread_block()、this_grid()、coalesced_threads()、this_cluster()。
partition 是 collective 操作，group 內所有 thread 都必須參與；若把 group 建立放在「不是所有 thread 都會到達」的條件分支中，會導致 deadlock 或 data corruption（建立 group 的程式碼必須讓全員都執行到）。

重疊是 opportunistic（機會性）、不保證真的並行；若把邏輯正確性建立在「一定並行」上可能導致 deadlock。重疊只是效能加分，正確性不能依賴它。
即使 primary 顯式 trigger，secondary 仍可能在 primary 結果寫回 global memory「之前」就 launch；少了 cudaGridDependencySynchronize（或等效手段），secondary 的相依工作會讀到尚未寫回的舊資料。

GC 是輕量 context，建立時即綁定一組特定 GPU 資源（目前是 SMs 與 work queues），targeting 它的工作只能用其 provisioned 資源；不需改 kernel，只改 host 端。
四步驟：取得資源（cudaDeviceGetDevResource）→ split SM 資源（cudaDevSmResourceSplit(ByCount)）→ 產生 descriptor（cudaDevResourceGenerateDesc）→ 建立 GC（cudaGreenCtxCreate）。不保證真正並行，整套技術只是「移除可能阻礙並行的因素」。

需 runtime ≥ 11.7「且」driver ≥ 515（AND 關係，缺一即整個程式回退 eager）；不需任何 compiler 支援，pre-11.7 編譯的 SASS/PTX 也能 lazy load（仍需 11.7+ runtime）。
含 managed variable 的 module 仍會 eager 載入。Error Log 用 CUDA_LOG_FILE（值可為 stdout / stderr / 合法檔案路徑）啟用。

init() 第二個參數是 expected arrival count，即「參與 thread 在某 thread 從 wait() 解除前會呼叫 arrive() 的總次數」。
arrive() 只回報「我到了」並使 countdown 遞減、不阻塞，呼叫後可立即去做獨立計算（split arrive/wait）；真正阻塞在 wait()。even phase parity = 0、odd = 1，初始 parity 為 0，有效值只有 0 / 1。

commit 會被 coalesce：warp converged 時序列只 +1、所有操作 batch 在同一 stage；fully diverged 時序列 +32、操作散到 32 個不同 stage，而每個 thread 的「感知序列」TL 仍小於「實際序列」PL。wait_prior 實際等到實際序列的 PL-N，當 TL < PL 時就會多等更新的 batch（極端時每個 thread 等全部 32 個）。
避免方式：讓 commit 由 converged threads 發出，使感知序列對齊實際序列；若 commit 前程式碼讓 threads diverge，應先用 __syncwarp 重新收斂再 commit。

CC 8.0+，唯一方向是 global → shared（資料直接落 shared、繞過 register）。複製 4/8 bytes 走 L1 ACCESS（資料同時快取於 L1），16 bytes 走 L1 BYPASS（不污染 L1）。
預設每個 thread 只等「自己」發出的 LDGSTS 複製；若用 LDGSTS prefetch 給其他 thread 用，在與完成機制同步後還必須額外加一個 __syncthreads()。

bulk-asynchronous copy 處理 1D 連續陣列、免 tensor map；bulk-tensor asynchronous copy 處理多維（最多 5D），需 host 端用 cuTensorMapEncodeTiled 建立的 tensor map。
讀用 shared memory barrier 的 transaction count（block 內任一 thread 可等）；寫用 bulk async-group（commit_group + wait_group，僅發起 thread 可等）。注意 bulk 與 non-bulk(LDGSTS) 的 async-group 彼此獨立、wait 指令不可混用。

唯一方向 register → distributed shared memory（不可反向、不碰 global memory），單次複製 4 / 8 / 16 bytes，需 CC 9.0+（依賴 cluster 機制）。
用 shared memory barrier（mbarrier）標示完成，而非回傳值或 fence；只透過低階 cuda::ptx::st_async（libcu++）暴露，無高階包裝。

在 Blackwell（compute capability 10.0）引入，結合 Fixed Work per Thread Block（load balancing / preemption）與 Fixed Number of Thread Blocks（reduced overheads）兩者的優點。核心動作：嘗試取消「尚未開始執行」的 block，成功則竊取其 index 來做它的工作（work stealing）。
UB 判定關鍵在於「兩個 request 之間有無插入 query 觀察」：在已 query 並斷定第一個失敗「之後」再發 request 即為 UB；連續發兩個 request、之後才 query 則合法。

hitRatio = 1.0 時硬體嘗試把整個 32KB 都快取進只有 16KB 的保留區，會不斷驅逐 cache line 來保留最近用到的 16KB（thrashing），持久化好處被抵銷。
hitRatio = 0.5 只隨機把 32KB 中的 16KB 標為 persisting，剛好放進 16KB 保留區、不抖動；hitRatio < 1.0 也可控制多個並行 stream 各自在 L2 快取多少、減少互相驅逐自己或對方的 persisting cache line。

根因是 cumulativity（累積性）：GPU 在執行當下無法分辨哪些 in-flight write 是 source level 真正保證要可見、哪些只是碰巧可見，fence/flush 只能保守地把「所有」in-flight 操作納入等待，因而不必要地等到通訊 kernel 較慢的 NVLink/PCIe 寫入。
domains 讓每個 kernel launch 取得一個 domain ID、writes 與 fences 都標記該 ID，fence 只 order 與其 domain 相符的 writes，把通訊 kernel 放到不同 domain 即可隔離流量。跨 domain 的 ordering 須用 system-scope fence（同 domain 內 device-scope 即足夠）。

device pointer / event handle 只在「建立它的 process」內有效，跨 process 無效；故須建立 process-portable handle（cudaIpcGetMemHandle），經 OS IPC（shared memory 或檔案）傳遞，對方再以 cudaIpcOpenMemHandle 還原成本地 pointer。
Legacy IPC 限制：僅 Linux、不支援 cudaMallocManaged、收發雙方須以相同 driver/runtime 編譯執行。cudaMalloc 可能從大區塊子分配，IPC 會分享「整個底層 block」而洩漏其他子分配，對策是只分享 2 MiB 對齊大小的 allocation。

cuMemMap 只建立 VA↔實體的映射、「不會」讓位址可存取；必須在來源與存取端 device 上呼叫 cuMemSetAccess（設 CU_MEM_ACCESS_FLAGS_PROT_READWRITE），否則 kernel 存取映射空間會 crash。
釋放順序固定為 cuMemUnmap → cuMemRelease → cuMemAddressFree（嚴格此序，顛倒會出錯）；OS-specific handle 另需 fclose 收尾，fabric 則不需。

用 OS 指派的 NUMA node id（numaID），「不等於」device ordinal；以 cuDeviceGetAttribute + CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID 取得（關聯最近的 host node）。
支援 cuMemCreate（VMM）與 cudaMemPoolCreate（Stream Ordered Memory Allocator）兩種 allocator。多節點須把 requestedHandleTypes 設為 CU_MEM_HANDLE_TYPE_FABRIC，並用 cuMemExportToShareableHandle / cuMemImportFromShareableHandle 跨節點共享。

x_array 在 parent 的 local memory，private 於該 thread、對 child 不可見；把 local（或 shared）指標當 child launch 參數是非法的，從 child grid 解參考為 undefined。傳給 child 的儲存須明確配置於 global heap（cudaMalloc / new，或 global scope 的 __device__ 變數）。
CDP2 已移除 cudaDeviceSynchronize，child 寫入「不保證」對 parent 可見；parent 退出前取回 child 修改的唯一方式是把一個 kernel 送進 cudaStreamTailLaunch stream。

生命週期：register → map → 取 device pointer/array（buffer 用 cudaGraphicsResourceGetMappedPointer、array 用 cudaGraphicsSubResourceGetMappedArray）→ 在 kernel 使用 → unmap → unregister；register / unregister 各一次，map / unmap 可任意多次。
register 是昂貴操作故每資源理想只做一次，但「每個打算使用該資源的 CUDA context 都必須各自註冊一次」，註冊結果不可跨 context 共用。

Linux fd 匯入後「CUDA 取得所有權」，匯入成功後再使用該 fd 屬 undefined behavior；Windows NT handle（OPAQUE_WIN32）CUDA「不」取得所有權，應用程式須在不需要時自行 CloseHandle，且該 NT handle 持有資源 reference，須先釋放才能釋放底層記憶體。
匯入 D3D12 資源（尤其 committed resource）一律須設 cudaExternalMemoryDedicated flag。

版本引數的語意是「要綁定哪個 typedef 的 ABI」；若改傳更高/動態的版本，部署到更新的 driver 時可能取回比 typedef 還新的符號（如未來的 _v3），ABI/簽章不符造成 undefined behavior，故須用「對應 typedef 的常數」硬編。
VERSION_NOT_SUFFICIENT = driver 中有此符號、但它比傳入的 cudaVersion 更晚加入（升高 cudaVersion 即可，driver 已夠新）；SYMBOL_NOT_FOUND = driver 中找不到符號（driver 太舊或名稱拼錯，此時 cudaVersion 給多少都無關）。

第四章「CUDA Features」橫跨記憶體、執行模型、互通與平台控制，可歸納為六條主線：

• Unified Memory：hardware-coherent（合併 page table、cache-line 粒度）vs software-coherent（各自 page table、fault+migration、page 粒度）決定 fault/contention 行為；concurrentManagedAccess = 0 平台 kernel 執行中 GPU 獨佔所有 managed data，須先同步或用 stream-attach 縮粒度；performance hints（prefetch / mem-advise）只影響效能不影響正確性。
• CUDA Graphs：definition → instantiation → execution，一次 instantiate 多次 launch；只改參數用 graph update（下次 launch 生效）、拓撲變才 re-instantiate；conditional（IF/WHILE/SWITCH）在 device 上評估；memory node 固定的是「VA」（圖內共用 VA、圖間 virtual aliasing 共用 physical）；device graph 三模式 fire-and-forget / tail / sibling。
• 記憶體配置與管理：stream-ordered allocator 避開 cudaMalloc/Free 的跨 stream 同步；VMM 把「VA 保留」與「實體配置」解耦、提供 allocation 粒度 peer 控制（map 後須 SetAccess、釋放 unmap→release→addressFree）；EGM 用 NUMA id 表達放置、多節點靠 FABRIC handle；IPC 交換 handle 非 pointer。
• 非同步協調與資料搬移：async barrier 把 arrive（不阻塞、回 token）/ wait（阻塞）分離、用 parity 追蹤 phase；pipeline 多階段緩衝、warp diverge 會 over-wait（先 __syncwarp）；三種 async copy——LDGSTS（global→shared, 8.0）、TMA（global↔shared 多維, 9.0, tensor map）、STAS（register→dsmem, 9.0）；cooperative groups partition 是 collective（全員須到）；PDL 與 CLC 都是 opportunistic，正確性不可依賴並行。
• 平台與細緻控制：green context 綁定 SM/WQ 降干擾（不保證並行）；L2 cache control 用 set-aside + hitRatio 避免 thrashing；memory sync domains 以 domain ID 縮小 fence 撒網（跨 domain 需 system-scope）；lazy loading 需 runtime ≥ 11.7 AND driver ≥ 515。
• 互通與 Driver 存取：graphics interop（register/map 生命週期、per-context 註冊、不支援 D3D12）；external interop（UUID/LUID/GPU id 配對、fd vs NT handle 所有權、dedicated flag）；driver entry point access 用精確版本取址，兩個失敗碼分辨「cudaVersion 太低」與「driver 太舊」。
  共同精神：多數功能在 host 端、不需改 kernel 即可調控 GPU work 與兩側效能；正確性要求（coherency 粒度、stream/event ordering、SetAccess、fence scope、handle 所有權、版本精確）絕不能為效能而妥協。