面試陷阱題 (Exam Traps)
Purpose
本檔收錄第一章 (Introduction to CUDA) 最容易答錯與最常被混淆的觀念，專供面試與測驗前快速複習使用。每個陷阱以可折疊 (fold) callout 呈現：先看標題自我作答，再展開核對「陷阱是什麼 / 為何容易混淆 / 正確觀念」，並附對應概念筆記連結深入回讀。建議搭配 速查表 一起服用。
GPU 基礎
Trap: 2003 與 2006 兩個年份別搞混
- 陷阱是什麼：把「GPU pipeline 可程式化」與「CUDA 問世」混為一談，或記錯年份。
- 為何容易混淆：兩件事都跟「GPU 變得更能執行通用程式」有關，發生時間又相近。
- 正確觀念：2003 年繪圖 pipeline 的部分階段變成完全可程式化（仍限繪圖用途）；2006 年 NVIDIA 推出 CUDA，讓任意運算工作負載不經繪圖 API 即可使用 GPU throughput。記法：2003 = 繪圖可程式化、2006 = CUDA 通用運算。
- 回讀：GPU 運算基礎
Trap: 「GPU 一定比 CPU 快」是過度簡化
- 陷阱是什麼：認為任何程式搬到 GPU 都會更快。
- 為何容易混淆：常聽到「GPU 加速」一詞，容易以為 GPU 全面優於 CPU。
- 正確觀念：GPU 只在高度平行工作負載、且於相近價格與功耗範圍下勝出（更高 throughput 與 bandwidth）；對低延遲序列工作，CPU 反而更適合。GPU 是以較低單執行緒效能換取更高總 throughput。此外 FPGA 同樣省電，但程式設計彈性遠不如 GPU。
- 回讀：GPU 運算基礎
執行模型與 SIMT
Trap: SIMT 不等於 SIMD
- 陷阱是什麼：把 CUDA 的 SIMT 直接當成 CPU 向量指令的 SIMD，認為「整個 warp 只能走一條控制流、資料寬度固定」。
- 為何容易混淆：兩者都是「一條指令作用於多筆資料」，名稱也只差一個字，硬體上 warp 確實同時執行同一條指令，看起來很像 SIMD lane。
- 正確觀念：SIMT (Single-Instruction Multiple-Threads) 允許 每個 thread 有自己的控制流路徑，分歧時用 masking 處理；且 沒有固定 data-width，programmer 寫的是 per-thread 程式碼。SIMD 則是單一控制流、固定 data-width、per-vector 運算。關鍵差異就兩點：per-thread 控制流 + 無固定資料寬度。
- 回讀：執行模型與 SIMT
Trap: block 之間沒有執行順序保證
- 陷阱是什麼：寫出「block 1 讀取 block 0 算好的結果」這類跨 block 相依，或假設 block 會依 blockIdx 由小到大依序執行。
- 為何容易混淆：在 CPU 思維裡迴圈是有序的；而且 grid 看起來像一個整齊的索引陣列，直覺以為 block 會照順序跑。
- 正確觀念：grid 可能含數百萬 block，但 GPU 只有數十到數百個 SM。每個 block 由單一 SM 執行、通常跑到完成，block 之間可以任意順序（平行或串行）執行、無排程順序保證，因此不可有跨 block 資料相依。這正是「任意大 grid 能在任意大小 GPU 上執行」的前提。需要跨 block 同步時，唯一受支援途徑是 cluster (CC 9.0+) 搭配 Cooperative Groups 與 distributed shared memory。
- 回讀：執行模型與 SIMT
Trap: warp divergence 與「thread 數要 32 的倍數」
- 陷阱是什麼：(1) 以為 if/else 分支「免費」；(2) 把 block 的 thread 數隨意設成例如 100、250，忽略 warp 大小。
- 為何容易混淆：程式碼編譯通過、結果也正確，看不出效能損失；warp 是硬體排程單位，原始碼層級不會主動提醒你。
- 正確觀念：block 內 threads 以 32 個為一組成一個 warp，warp 內所有 thread 同時執行同一條指令。當同 warp 內 threads 走不同分支即 warp divergence，未走該分支的 lane 被 mask off 閒置，利用率下降；warp 內走相同控制流時利用率最高。thread 數雖可任意，但最好是 32 的倍數，否則最後一個 warp 會有部分 lane 全程閒置，導致 functional units 與記憶體存取次佳 (suboptimal)。
- 補充：不要依賴「warp 如何對應實體硬體」寫程式；違反 programming model 會導致 undefined behavior。
- 回讀：執行模型與 SIMT
Tile 程式設計
Trap: tile 是不可變的，且維度須為 2 的次方且編譯期已知
- 陷阱是什麼：想「就地修改」一個 tile、把 tile 當 kernel 參數傳遞，或用執行期才算出的變數（甚至非 2 的次方）當 tile 維度。
- 為何容易混淆：tile 看起來就像一個多維陣列/矩陣，很自然會以操作 array 的方式對待它。
- 正確觀念：tile 是 immutable——每個運算都產生新 tile，不會修改既有 tile；每個維度必須是 2 的次方且編譯期已知（值要在 kernel 執行前可決定，而非執行期才算出）；tile 不能當 kernel 參數，只能在 tile code 內建立與消耗；tile 也不一定有記憶體表示（由 compiler 決定放 registers / shared memory / 其他 SM 資源）。
- 回讀：Tile 程式設計
Trap: array (可變) vs tile (不可變) 性質幾乎相反
- 陷阱是什麼：把 array 與 tile 當成同一種東西，混用兩者的可變性與生命週期假設。
- 為何容易混淆：兩者都是「多維資料」，在 tile kernel 裡又經常一起出現（load 把 array 讀成 tile、store 把 tile 寫回 array）。
- 正確觀念：兩者性質幾乎相反——array：存於 device memory、可變 (mutable)、可當 kernel 參數、一定有記憶體表示、有 shape 與 dtype；tile：只存在於 tile code、block 區域、不可變 (immutable)、維度為 2 的次方且編譯期已知、不一定有記憶體表示、不可當 kernel 參數。記法：block 是執行單位、tile 是資料單位；一個 block 可建立多個不同 shape 的 tile。
- 額外易考點：load 與 store 的越界行為不對稱——load 越界依指定方式處理（如補零），store 越界寫入則被靜默丟棄。
- 回讀：Tile 程式設計
記憶體階層
Trap: global memory 並非「系統內到處可存取」
- 陷阱是什麼：看到 "global" 就以為它是「全系統共用、CPU/GPU 各處都能直接存取」的記憶體。
- 為何容易混淆："global" 這個字面意義太強，加上 unified virtual memory space 讓人誤以為位址通用就代表存取通用。
- 正確觀念：global memory 之名僅表示「該 GPU 內所有 SM 皆可存取」連接到該 GPU 的 DRAM，不代表整個系統各處皆可存取。CPU 端的 DRAM 另稱 system memory / host memory。雖然 CPU 與所有 GPU 共用單一 unified virtual memory space（每個位址範圍唯一不重疊、可由位址判斷屬於哪塊記憶體），但「位址統一」不等於「存取統一」。
- 回讀：GPU 記憶體階層
Trap: unified memory 仍要盡量減少搬移
- 陷阱是什麼：以為用了 unified memory 就「不用管資料搬移、不用管 locality，效能自動最佳」。
- 為何容易混淆：unified memory 讓單一配置可從 CPU 或 GPU 存取、由 runtime/硬體在需要時自動搬移，使用上「看起來」不需手動 cudaMemcpy。
- 正確觀念：unified memory 只是讓存取正確且省去手動複製，但最佳效能仍來自盡量減少搬移，並盡可能從直連該記憶體的處理器直接存取。自動搬移本身有成本，不是免費的。
- 另一個常見誤解 — mapped memory：它是可被 GPU 直接存取的 CPU memory，但走 PCIe / NVLINK，延遲高、頻寬低且無法用 parallelism 隱藏，並非 unified memory 或正確擺放資料的高效替代方案。
- 回讀：GPU 記憶體階層
Trap: register 是 per-thread、shared memory 是 per-block
- 陷阱是什麼：把 register 與 shared memory 的配置粒度搞反，或忽略 register 不足會導致 kernel 根本無法啟動。
- 為何容易混淆：兩者都是 SM 內極快的 on-chip 記憶體，常被一起提到。
- 正確觀念：register file 是 per-thread（存 thread 區域變數、由 compiler 配置）；shared memory 是 per-block（整個 thread block / cluster 共用，用於 threads 間交換資料）。排程條件：每 thread 所需 register 數 × block 內 thread 數 ≤ SM 可用 register；若 block 所需 register 超過 register file 大小，kernel 無法啟動，必須減少 block 內 thread 數。
- 回讀：GPU 記憶體階層
CUDA 平台與相容性
Trap: cubin 跨 major CC 不相容，僅同 major 內 minor 向前相容
- 陷阱是什麼：以為「為舊 CC 編的 cubin 一定能在更新的 GPU 上跑」，把 PTX 的向前相容套用到 cubin。
- 為何容易混淆：直覺認為「版本號越新越向下相容」；PTX 確實能 JIT 到更高 CC，容易把這個性質誤植到 cubin。
- 正確觀念：cubin 對應特定 SM 版本。規則建立在 CC 的 major/minor 上：同一 major 內，minor 大於或等於目標的 GPU 可載入該 cubin；跨 major 不相容。例：
sm_86(CC 8.6) 可在 CC 8.6、8.9 執行，但不能在 CC 8.0（minor 較低）或 CC 9.0（跨 major）執行。要兼顧未來 GPU 應改靠 PTX：儲存的 PTX 可在執行期 JIT 到相同或更高 CC（forward compatibility），但不能降到更低 CC。binary compatibility 也僅對 NVIDIA 工具產生、未經修改的 binary 保證。 - 回讀：CUDA 平台
Trap: compute_XY (PTX) vs sm_XY (cubin) 命名混淆
- 陷阱是什麼：把
compute_80與sm_80當同義詞混用。 - 為何容易混淆：兩者格式相近、都對應 Compute Capability，常一起出現在編譯指令中。
- 正確觀念：
compute_XY指 PTX（虛擬 ISA / IR，對應 CC，是可 JIT 的中介表示）；sm_XY指 cubin（針對特定 SM 版本的實體 GPU 二進位）。編譯流程：高階語言 → PTX (compute_XY) → cubin (sm_XY)；fatbin 則是可同時容納多個 target cubin 與 PTX 的容器，執行時挑最適 binary。 - 回讀：CUDA 平台
第二章：Programming GPUs in CUDA
Purpose
本區延續上方格式，收錄第二章 (Programming GPUs in CUDA) 最容易答錯與最常被混淆的觀念，涵蓋 CUDA C++/Python 的 host 端流程、SIMT 與 Tile 兩種 kernel 模型、非同步執行、Unified/System Memory 與 NVCC。用法相同：先看標題自我作答，再展開核對「陷阱是什麼 / 為何容易混淆 / 正確觀念」，並沿 回讀 連結深入。
CUDA C++
Trap: kernel launch 對 host 非同步、但 cudaMemcpy 是同步的
- 陷阱是什麼：以為
<<<>>>啟動後 kernel 已跑完、可直接讀輸出；或反過來以為cudaMemcpy也像 launch 一樣立即返回。 - 為何容易混淆：兩者都是「把工作丟給 GPU」，直覺以為同步性一致。
- 正確觀念：kernel launch 對 host thread 非同步，立即返回、host 會在 kernel 完成（甚至開始）前就繼續往下跑；要讀結果或
cudaFree前必須同步（最簡單cudaDeviceSynchronize，多 stream 時改用cudaStreamSynchronize/ events 以免過度同步）。反觀cudaMemcpy是 同步的：複製完成前不返回、會阻塞 host thread。要非同步搬移得用cudaMemcpyAsync搭配 pinned memory。記法：launch 非同步、memcpy 同步。 - 回讀：同步與完整流程
Trap:
<<<>>> 不回傳 cudaError_t，且查到 cudaSuccess 不代表 kernel 成功
- 陷阱是什麼：想用
cudaError_t err = kernel<<<>>>(...)接 launch 錯誤；或 launch 後立刻cudaGetLastError()拿到cudaSuccess就認定 kernel 已正確跑完。 - 為何容易混淆：其他 runtime API 都回傳
cudaError_t，自然以為 launch 也是；而且 launch 非同步，當下根本還沒執行。 - 正確觀念：triple chevron 是語言擴充、不回傳
cudaError_t，要靠cudaGetLastError（取回並重設）/cudaPeekAtLastError（不重設）查 per-host-thread 的錯誤狀態。launch 後立即檢查只能抓「launch 參數 / execution configuration 錯誤」與殘留的非同步錯誤；cudaSuccess不代表 kernel 已執行或成功。要抓 kernel 執行期間的非法存取等非同步錯誤，需在cudaDeviceSynchronize()之後再檢查。非同步錯誤會被之後每個 API 持續回傳，直到cudaGetLastError清除。 - 回讀：錯誤檢查與修飾符
Trap: explicit management 下 host pointer 不能直接傳給 kernel
- 陷阱是什麼：用
cudaMalloc配了 device buffer，卻把 host 的A（malloc/cudaMallocHost指標）直接傳進 kernel；或在 host 端 dereference device 指標。 - 為何容易混淆：unified memory（
cudaMallocManaged）用單一指標兩端共用，讓人以為 explicit 也行；指標型別又長得一模一樣。 - 正確觀念：explicit memory management 下 host pointer 與 device pointer 在分離的位址空間，不可互相直接 dereference。kernel 必須收 device 指標（
devA…），輸入要先cudaMemcpyH2D，算完再 D2H 拷回 host buffer 才能在 CPU 讀。只有 unified memory 才是「單一指標、兩端皆可存取、driver 自動搬移」。 - 回讀：CUDA C++ 記憶體管理
CUDA Python
Trap: CuPy 不會隱式搬移；ndarray 只在 host 或 device 其一
- 陷阱是什麼：把 Numpy（host）array 直接傳給
@cuda.jitkernel，或把 CuPy（device）array 傳給一般 Python 函式，期待自動搬到對的位置。 - 為何容易混淆：Python 平常不在意資料位置，CuPy 介面又與 Numpy 幾乎一樣。
- 正確觀念：ndarray 只存在於 host 或 GPU 其一，不會兩邊都有；CuPy 不會隱式在 CPU↔GPU 複製（複製昂貴、過度複製傷效能），要程式設計師明確控制：H→D 用
cp.array(host)、D→H 用cp.asnumpy(device)。host array 傳 kernel → error；GPU array 傳一般 Python 函式 → error。另記：CuPy 未指定 dtype 時預設 float32；kernel 內用 ndarray 會攜帶 extent，自動邊界檢查。 - 回讀：CUDA Python 入門
Trap: CUDA Python 錯誤是 exception（非回傳碼），但 launch 一樣非同步
- 陷阱是什麼：照 C++ 思維去接
cudaError_t回傳值；或以為 Python 的 kernel launch 是同步的。 - 為何容易混淆：API 概念與 C++ 一致，容易把錯誤處理模型也照搬。
- 正確觀念：在 Python，CUDA 錯誤會 拋 exception（未捕捉會終止程式），用
try/except攔截——例如block_size=2048（> 1024 上限）會得CUDA_ERROR_INVALID_VALUE。同時 kernel launch 對 host 仍是非同步，需 device-wide 同步（cp.synchronize()/cuda.synchronize()/device.synchronize()，等同cudaDeviceSynchronize）；不過cp.asnumpy這類 D→H 複製會隱式等待 kernel 完成。 - 回讀：CUDA Python 入門
SIMT Kernels
Trap:
__syncthreads() 必須整個 block 都到達，否則 deadlock
- 陷阱是什麼：把
__syncthreads()放進只有部分 thread 會進入的 divergent 分支，如if (threadIdx.x < N) { ... __syncthreads(); }。 - 為何容易混淆：把它當成普通的「等一下」呼叫，沒意識到它是 block 層級 barrier，要求全員到齊。
- 正確觀念：
__syncthreads()（Pythoncuda.syncthreads()）是 block barrier：block 內所有 thread 都到達前，沒有任何 thread 能往下走，並保證 barrier 前對 shared memory 的寫入對 barrier 後可見。若有 thread 永遠到不了（被分支擋掉），就 deadlock / 未定義行為。需要條件分支時，把 barrier 放在所有 thread 都會執行到的位置。它只同步單一 block，跨 block 要靠 thread block clusters / Cooperative Groups / atomics。 - 回讀：SIMT 基礎與 Thread 階層
Trap: local memory 其實在 device memory，不是 SM 上
- 陷阱是什麼：看到 "local" 就以為它在 SM 上，是像 register / shared 一樣快的 thread 私有快取。
- 為何容易混淆："local" 字面像「近、快、本地」，又和 register 一樣是 thread scope、由 compiler 管理。
- 正確觀念：「local」指的是邏輯 scope（thread local），不是實體位置；local memory 實體在 off-chip 的 device（global）memory space，延遲與頻寬等同 global memory，同樣受 coalescing 規範。它是 register spilling、無法以常數索引的陣列、過大 struct 的去處。記法：register / shared 在 SM；global / constant / local / L2 在 device。
- 回讀：SIMT 裝置記憶體空間
Trap: global memory coalescing 不要求「連續」，permuted 也算合併
- 陷阱是什麼：以為只有「連續 thread 取連續位址」才算 coalesced；或反過來以為隨便存取硬體都會救回來。
- 為何容易混淆：教學常以「連續存取」當範例，讓人把「連續」當成 coalescing 的定義。
- 正確觀念：global memory 以 32-byte transaction 為單位，warp 把 32 個 thread 的請求合併成所需的最少 transaction。真正條件是 warp 內 thread 存取同一批 32-byte segment——連續或 permuted（排列）都算合併，目標是最大化 used bytes / transferred bytes。完美合併（連續 4-byte word）= 128 bytes / 4 個 transaction / 100%；病態（stride ≥ 32 bytes，如轉置寫
c時 stride = ld）= 每 thread 一個 transaction、1024 bytes 流量只用 128 bytes = 12.5%。 - 回讀：SIMT 記憶體效能
整欄 = 32-way）
- 陷阱是什麼：宣告
__shared__ float s[32][32]就以為 shared memory 一定快，忽略 bank conflict。 - 為何容易混淆：shared memory 在 SM 上、延遲低，容易以為任何存取樣式都高效。
- 正確觀念：shared memory 分 32 個 bank，連續 32-bit word 對映連續 bank（
word % 32）。同 warp 多 thread 存取同一 bank 的不同位置 → 序列化（bank conflict）。s[32][32]按「整欄」存取（s[threadIdx.x][threadIdx.y]，threadIdx.x 當第一索引）stride = 32 → 全落 bank 0 → 32-way conflict；按「整列」stride = 1 → 無衝突。解法是 padding：欄維 +1 宣告成[32][33]，列長 33 與 32 互質，使整欄存取錯開到 32 個不同 bank。例外（無衝突）：同 warp 存取同一位置時，讀會 broadcast、寫只有一個 thread 寫入（哪個未定義）。 - 回讀：SIMT 記憶體效能
Trap: 非 atomic 的 read-modify-write 會有 race condition
- 陷阱是什麼：多個 thread 對同一 global 位置做
s[0] = s[0] + x，期待得到正確總和。 - 為何容易混淆：單執行緒看起來沒問題；小資料量偶爾還「碰巧」對。
- 正確觀念：grid 層級沒有內建全域同步，跨 thread 更新同一位置必須用 atomic。非 atomic 的 RMW 會遺失更新、結果偏小，且每次執行、在不同 SM 數的 GPU 上都可能不同。用
atomicAdd/cuda::atomic_ref（C++，可指定 thread scope，如thread_scope_device）或cuda.atomic.add（Python）。atomic 會強制同步、有成本，應節制使用：先在 shared memory 局部歸約，再由單一 thread 做一次 atomic。 - 回讀：Atomics/Cooperative/Occupancy
Trap: occupancy 太低——block 太小反而吃虧
- 陷阱是什麼：以為 block 開很小（如 32 threads）或一味壓低資源就能塞滿 SM、拉高 occupancy。
- 為何容易混淆：直覺認為只要 thread 總量沒超過上限就沒事，忽略「每 SM 最多 block 數」這個獨立上限。
- 正確觀念：occupancy = active warps ÷ SM 最大 active warps，越高越能隱藏延遲。限制資源三類：threads/block、shared memory/block、registers/thread，各有 per-SM 與 per-block 上限。block 太小會先撞到
maxBlocksPerMultiProcessor（如 CC 10.0 為 32）：<<<512, 32>>>→ 32 blocks × 32 = 1024 / 2048 = 50%；<<<512, 768>>>受 2048 thread 限 → 每 SM 2 blocks → 75%。nvcc --maxrregcount壓 register 可能塞更多 block，但過低會 spill 到 local memory，需實測權衡。block 排到哪個 SM 不可控、不可查、無順序保證。 - 回讀：Atomics/Cooperative/Occupancy
Tile Kernels
Trap: tile 不可變、維度須 2 的次方且編譯期已知；C++ 第二 chevron 參數恆為 1
- 陷阱是什麼：想就地修改一個 tile、用 runtime 才算出的值（甚至非 2 次方）當 tile 維度，或在 tile launch 把
<<<grid, 256>>>的第二參數當成 thread 數來填。 - 為何容易混淆：tile 看起來就像多維陣列；tile kernel 又沿用 SIMT 的 triple chevron 語法。
- 正確觀念：tile 的 shape 與 dtype 必須編譯期已知、每維為 2 的次方，具 value semantics（複製產生獨立副本、但成本低），programmer 不配置/釋放其記憶體。C++ tile kernel 的 triple chevron 第二參數必須是
1（thread 數由 compiler 決定），寫其他值即錯；runtime 維度要用ct::extents的 brace 形式（_ic標編譯期、普通變數標 runtime）。tile-space 邊界：partial OOB 的 unmasked load 是 UB，masked 變體才安全；完全落在 array 外的 tile 連 masked 也救不了。 - 回讀：Tile Kernel 結構
Trap: tile atomic「整個 call 非 atomic」；C++ 預設 system scope vs Python 預設 device scope
- 陷阱是什麼：以為對一個 tile 做
atomic_add是把整塊當一次不可分割的操作；或假設 C++ 與 Python 的 atomic 預設 thread scope 相同。 - 為何容易混淆：一行呼叫看起來像單一原子操作；兩語言 API 名稱又相近。
- 正確觀念：tile atomic 對 tile 的每個元素各做一次 atomic，整個 call 並非 atomic，元素間先後順序未指定。thread scope 預設兩語言不同：C++ 省略時是 system-wide、Python 預設是 device。cross-block 合併需 device scope；intra-block 競爭用 block scope 即可（加法可交換，故未指定的順序不影響結果）。要把 tile 加總成 scalar 應改用 tile reduction，別用 intra-block atomics 硬湊；Python 不需舊值時用
TiledView.atomic_add（lower 成 PTX atomic reduction，較快）。 - 回讀：Tile Atomics 與最佳化
Trap:
__restrict__ 用錯是 undefined behavior
- 陷阱是什麼：為了效能對所有 pointer 都加
__restrict__，包括實際上可能 alias（指向重疊區域）的 pointer。 - 為何容易混淆：加了通常會更快又能編譯通過，看不出問題。
- 正確觀念：
__restrict__是對 compiler 的保證——該記憶體區域在此 pointer 生命週期內只透過這個 pointer 存取。成立時 compiler 不必為 overlap 做保守同步，load/store 可交錯、pipeline（這是 tile C++ 良好記憶體效能的關鍵）。但若該區域其實可被另一個 pointer 存取，仍標__restrict__就是 undefined behavior。相關效能標註：ct::assume_aligned(p, 16_ic)（16-byte 對齊，runtime 須真對齊否則 UB，是partition_view走 TMA 的前提；cudaMalloc保證 ≥ 16-byte）。 - 回讀：Tile Atomics 與最佳化
Asynchronous Execution
Trap:
cudaMemcpyAsync 用非 pinned host memory 會悄悄退化成同步
- 陷阱是什麼：把
cudaMemcpy換成cudaMemcpyAsync就以為傳輸一定會與運算重疊。 - 為何容易混淆：API 名字有 "Async"，呼叫也立即返回，看起來就是非同步。
- 正確觀念：要真正非同步、能與其他工作重疊，host buffer 必須是 pinned / page-locked（
cudaMallocHost/cudaHostAlloc）。用 pageable 記憶體時cudaMemcpyAsync仍正確，但會退化成同步、無法重疊，喪失效能優勢。重疊還需把工作放到不同 stream，並用 events /cudaStreamSynchronize做細粒度同步，而非cudaDeviceSynchronize（會等所有 stream、過度同步、抹掉重疊效益）。 - 回讀：非同步 Streams 與 Events
Trap: default stream 是 blocking，會把「看似能並行」的 kernel 串起來
- 陷阱是什麼：把幾個 kernel 丟到不同 stream，但中間夾了一個沒指定 stream 的操作，卻假設它們會並行。
- 為何容易混淆：每個 kernel 都丟進「自己的」stream，直覺以為彼此獨立。
- 正確觀念：未指定 stream 的操作進入 legacy default stream（NULL stream / ID 0），它是 blocking 且所有 host thread 共享：排入時會與所有其他 blocking stream 同步（等它們完成、也擋住後續）。要避免就用
cudaStreamCreateWithFlags(..., cudaStreamNonBlocking)的 non-blocking stream，或啟用 per-thread default stream（--default-stream per-thread或CUDA_API_PER_THREAD_DEFAULT_STREAM）。另外 stream 內是 in-order，跨 stream 無順序保證，要靠 events /cudaStreamWaitEvent；stream callback / host 函式內不可呼叫任何 CUDA API。 - 回讀：Callbacks/排序/Graphs
Unified/System Memory
Trap: unified memory 仍要盡量減少搬移
- 陷阱是什麼：以為用了
cudaMallocManaged就不用管資料位置與 locality，效能自動最佳。 - 為何容易混淆：單一指標、省去手寫
cudaMemcpy，使用上「看起來」沒有搬移。 - 正確觀念：unified memory 只保證存取正確並省去手動複製，driver / 硬體仍會在背後按需遷移（software coherence = page 粒度、hardware/ATS = cache line 粒度），自動搬移有成本、不是免費。最佳效能仍來自減少搬移、讓資料常駐在直連的 processor，並用
cudaMemAdvise（放置提示）與cudaMemPrefetchAsync（kernel 前預先搬、與運算重疊）調效。注意 limited support（Windows / WSL / Tegra）下 GPU 活動期間 CPU 不可存取 managed memory、且不允許 oversubscription。 - 回讀：Unified 與 System Memory
Trap: mapped (zero-copy) memory 是高延遲，不是高效替代品
- 陷阱是什麼：把 mapped host memory 當成 unified memory 或正確擺放資料的「免搬移高效版」，用來滿足 kernel 大部分記憶體需求。
- 為何容易混淆：kernel 能用同一個 host 指標直接存取、不必
cudaMemcpy，看起來很方便。 - 正確觀念：mapped memory 資料一直留在 CPU 記憶體，kernel 每次存取都要走 PCIe / NVLink 交易，延遲高、頻寬低，難以用 parallelism 隱藏；它一律是 page-locked，而對它的 atomic 從 host / 其他 GPU 看不是 atomic。
cudaMallocHost/cudaHostAlloc(mapped) 回傳的 host 指標可直接用；但cudaHostRegister註冊既有配置後，kernel 必須改用cudaHostGetDevicePointer取得的 device 指標，不能用原 host 指標。對比 unified memory：首次遷移後可享完整 GPU 記憶體頻寬。 - 回讀：Unified 與 System Memory
NVCC
Trap: cubin 跨 major CC 不相容——別把 PTX 的向前相容套到 cubin
- 陷阱是什麼：以為為舊架構編的
sm_XYcubin 一定能在更新的 GPU 上跑。 - 為何容易混淆：直覺「版本越新越向下相容」；PTX 確實能 JIT 到更高 CC，容易誤植到 cubin。
- 正確觀念：
compute_XY是 virtual ISA → 產生 PTX（可 JIT、向前相容）；sm_XY是 real hardware ISA → 產生 cubin，以 SM 版本識別。cubin 規則：同一 major 內、minor ≥ 目標才能載入，跨 major 不相容（sm_86可在 CC 8.6 / 8.9，但不能在 8.0 或 9.0）。nvcc 預設為支援的最早/最低架構生成以求最大相容；device code 鏈為 C/C++ → PTX →ptxas→ cubin，可用 fatbin 同時內嵌多個 PTX / cubin。 - 回讀：NVCC 編譯器
Trap: PTX 只能 JIT 到「相同或更高」CC；
-arch=native 沒 PTX 也就沒向前相容
- 陷阱是什麼：以為內嵌 PTX 就能跑在任何 GPU（包含更舊的）；或圖方便用
-arch=native，卻期待產物能在別台更新的 GPU 上跑。 - 為何容易混淆：PTX 被說成「向前相容的中介表示」，容易過度推論成「萬用」。
- 正確觀念：儲存的 PTX 在執行期可 JIT 到相同或更高 CC（forward compatibility），但不能降到更低 CC。
-arch=sm_XY會同時內嵌 PTX 保留向前相容；但-arch=native只為當前 GPU 產 cubin、不含 PTX，因此無向前相容；-arch=compute_XY則是 PTX-only。補充易考點：自 CUDA 13 起__global__與__managed__/__device__/__constant__變數預設 internal linkage，跨單元引用（separate compilation，-rdc=true/-dc）時需特別注意。 - 回讀：NVCC 編譯器
第三章：Advanced CUDA
Purpose
本區延續上方格式，收錄第三章 (Advanced CUDA) 最容易答錯與最常被混淆的觀念，涵蓋進階啟動 / clusters、進階 streams 與相依啟動、批次傳輸與環境變數、PTX 與硬體模型、thread scopes 與 scoped atomics、非同步 barriers / pipelines、非同步資料複製與 L1/shared 配置、driver API、多 GPU 與功能導覽。用法相同：先看標題自我作答，再展開核對「陷阱是什麼 / 為何容易混淆 / 正確觀念」，並沿 回讀 連結深入。
進階啟動與 Clusters
Trap:
cudaLaunchKernelEx vs triple chevron 的使用時機
- 陷阱是什麼：以為
<<<>>>能表達所有啟動設定；或反過來凡事都改寫成囉嗦的cudaLaunchKernelEx。 - 為何容易混淆：兩者都只是「啟動 kernel」，
cudaLaunchKernelEx看起來只是更冗長的同義寫法。 - 正確觀念：triple chevron 只能表達四個固定參數——grid 維度、block 維度、dynamic shared memory（預設 0）、stream（預設 default stream）。一旦需要附加 launch attribute（如 cluster 維度
cudaLaunchAttributeClusterDimension、L1/shared carveoutcudaLaunchAttributePreferredSharedMemoryCarveout、PDL 的 stream serialization），就必須改用cudaLaunchKernelEx+cudaLaunchConfig_t（含attrs指標與numAttrs），且不需修改 kernel 原始碼即可逐次啟動附加任意數量屬性。無額外屬性時<<<>>>已足夠，不必為了用而用。 - 回讀：進階啟動與 Clusters
Trap: cluster 與 grid 的關係、可移植上限 8 blocks、Blocks-as-Clusters 佔位符
- 陷阱是什麼：以為啟動 cluster 後 grid 要改以 cluster 數枚舉、cluster 想開多大都行；或把
__block_size__啟用 Blocks-as-Clusters 後的第二個<<<>>>參數當 thread/block 數來填。 - 為何容易混淆：cluster 是「block 之上」的新層級，直覺以為 grid 單位也跟著改；各種維度約束細節又容易記錯。
- 正確觀念：用
cudaLaunchAttributeClusterDimension或__cluster_dims__時，grid 仍以 thread block 數枚舉，且 grid 各維度必須可被 cluster 各維度整除。可移植 (portable) 的 cluster 大小上限為 8 個 thread blocks；超過 8 需 opt-in 查詢、不可移植。只有__block_size__帶第二個 tuple（Blocks as Clusters）時<<<>>>第一參數才變成 cluster 數；此模式若要帶 dynamic shared memory / stream，第二個<<<>>>參數必須是佔位符1，填其他值是 undefined behavior。Thread block clusters 需 CC 9.0+（同 cluster 的 blocks 保證同時在單一 GPC 執行）。 - 回讀：進階啟動與 Clusters
進階 Streams 與相依啟動
Trap: stream priority 只是 hint、不搶佔；NULL stream / 隱式同步會阻斷並行
- 陷阱是什麼：用 stream priority 來「強制」工作順序，或把幾個 stream 設不同優先級就以為高優先一定先做完；又或在不同 stream 指令之間插入 NULL stream 指令 /
cudaMalloc卻假設仍能並行。 - 為何容易混淆：priority 一詞讓人以為是強制排序；NULL stream 看起來只是「沒指定 stream」很無害。
- 正確觀念：stream priority 只影響 scheduler 挑選待處理任務的順序，是 hint 而非保證，不會搶佔 (preempt) 已在執行的工作，也不會中途重評估 work queue（範圍由
cudaDeviceGetStreamPriorityRange()回[greatest, least]）。需要嚴格順序要用 event /cudaStreamWaitEvent。此外，host 在兩個不同 stream 指令之間若發出對 NULL stream 的任何指令、device/pinned 配置、memset、L1↔shared 切換等，會隱式序列化跨 stream 工作；用cudaStreamNonBlocking的 non-blocking stream 可避免 NULL stream 阻斷。準則：獨立操作先發、同步儘量延後。 - 回讀：進階 Streams 與相依啟動
Trap: PDL 三要件缺一不可；未 record 的 event 永遠回傳 success
- 陷阱是什麼：只在 secondary kernel 加
cudaGridDependencySynchronize()、或只設 launch attribute，就以為啟用了相依啟動重疊；或查詢一個還沒cudaEventRecord的 event 得到 success 就以為「等到了」。 - 為何容易混淆：PDL 三個元件分散在 primary kernel、secondary kernel、launch 端；event 查詢回 success 看起來就是完成了。
- 正確觀念：Programmatic Dependent Launch (PDL) 三要件缺一不可——primary 端
cudaTriggerProgrammaticLaunchCompletion()、secondary 端cudaGridDependencySynchronize()、secondary 用cudaLaunchKernelEx帶cudaLaunchAttributeProgrammaticStreamSerialization（programmaticStreamSerializationAllowed = 1）。即使三者齊全，重疊也只是「機會」非「保證」，取決於硬體資源與兩個 kernel「何時產出 / 何時消費」相依資料的結構。另外：未被 record 的 event 在任何 wait/query 永遠回傳 success，必須自行確保cudaEventRecord已先呼叫，否則會「以為等到了其實沒等」。只表達相依、不計時的 event 應用cudaEventDisableTiming建立以提升效能。 - 回讀：進階 Streams 與相依啟動
批次傳輸與環境變數
Trap: ephemeral pointer 必須用 DuringApiCall；環境變數要在啟動 app 前設定
- 陷阱是什麼：把 stack buffer（ephemeral pointer）丟進
cudaMemcpyBatchAsync卻用cudaMemcpySrcAccessOrderStream；或在程式內才setenv("CUDA_MODULE_LOADING", ...)後期待生效。 - 為何容易混淆：batch copy 的 attribute 種類多、語意細；環境變數在程式內設看起來理所當然會被讀到。
- 正確觀念：
cudaMemcpyBatchAsync的srcAccessOrder必須對應記憶體種類——pinned/managed 用...Stream（會 block 到先前 kernel 完成）、stack 上的 ephemeral pointer 必須用...DuringApiCall（否則 async copy 真正執行前指標可能失效）、heapmalloc且無 hardware managed/coherent access 用...Any。attrsIdxs[i]是「第 i 個 attribute 套用的第一個 transfer index」，failIdx傳nullptr是安全的。CUDA 環境變數（CUDA_DEVICE_MAX_CONNECTIONS減少 false dependency 序列化、CUDA_MODULE_LOADING=EAGER把載入移到初始化、CUDA_VISIBLE_DEVICES、CUDA_LAUNCH_BLOCKING）應在啟動 application 前設定；在 application 內設定可能完全無效。 - 回讀：批次傳輸與環境變數
使用 PTX 與硬體模型
Trap: independent thread scheduling 後不可再假設 warp lockstep（需
__syncwarp，CC 7.0+）
- 陷阱是什麼：沿用 CC 7.0 以前的 warp-synchronous code（如免同步的 intra-warp reduction），假設同 warp threads 每條指令都 lockstep，不加任何同步就讀彼此寫入的資料。
- 為何容易混淆：warp 一次仍執行一條共同指令，「看起來」像永遠 lockstep；舊 code 在 Volta 之前能正確跑，移植後又常碰巧不報錯。
- 正確觀念：自 CC 7.0 (Volta) 起引入 independent thread scheduling，每個 thread 有獨立的 program counter 與 call stack，可在 sub-warp 粒度 diverge / reconverge，不再保證 warp lockstep，參與某條指令的 thread 集合可能與預期不同。任何依賴隱式 warp-synchronous 行為的 code 都應重新檢視，並以
__syncwarp()顯式同步才跨世代正確。注意__syncwarp()同步的是 warp、__syncthreads()同步的是 block，兩者別搞混。 - 回讀：使用 PTX 與硬體模型
Trap: async proxy 需 proxy fence；inline PTX 需
volatile
- 陷阱是什麼：TMA / tensor core 等非同步路徑後，直接用 normal load/store 讀寫同一位址而不下 proxy fence；或內嵌 PTX 沒加
volatile，被編譯器最佳化掉或重排。 - 為何容易混淆：LDGSTS 那種 generic proxy 對「先前」的 normal access 仍保證順序，容易誤以為 async proxy 也一樣。
- 正確觀念：normal load/store 走 generic proxy；
LDGSTS/STAS/REDAS是 async thread on generic proxy；TMA /tcgen05.*/wgmma.mma_async.*走 async proxy。對 async proxy，先前與後續的 normal access 都不保證順序，跨 proxy 必須用 proxy fence 才能正確排序（generic proxy 至少保證「先前」的同址 access 順序）。手寫 inline PTX 要用asm volatile防止編譯器最佳化掉或重排該段。手寫 PTX 是最後手段，一般優先用既有 intrinsic 或cuda::ptx。 - 回讀：使用 PTX 與硬體模型
Thread Scopes 與 Scoped Atomics
Trap: thread scope 選太寬會慢、選太窄會錯（coherency point）
- 陷阱是什麼：為了求快一律用
thread_scope_block，或保險起見一律用thread_scope_system。 - 為何容易混淆：「scope 越窄越快」是真的，容易過度推論成「都選最窄」；system 最安全，也容易過度推論成「都選最寬」。
- 正確觀念：thread scope 對應記憶體階層的 point of coherency（block=L1、cluster/device=L2、system=L2 + connected caches），由窄到寬 thread → block(.cta) → cluster(.cluster) → device(.gpu) → system(.sys)。scope 越窄、coherency point 越靠近核心、同步成本越低；但 scope 必須涵蓋所有需要互相觀察 / 同步的 threads——兩個不在同一 block 的 thread 卻用
thread_scope_block，行為不正確。原則是「選能滿足正確性的最窄 scope」，先正確再求快（另記：shared memory atomics 比 global memory atomics 快）。 - 回讀：Thread Scopes 與 Scoped Atomics
Trap:
memory_order_relaxed 雖快，但不保證跨執行緒可見順序
- 陷阱是什麼：producer-consumer（先寫
data，再 setreadyflag）兩端都用memory_order_relaxed，以為是 atomic 就安全。 - 為何容易混淆：relaxed 確實是 atomic、也最快，名字只是「寬鬆」沒明示「會出錯」。
- 正確觀念：
memory_order_relaxed只保證該 atomic 本身的原子性，不提供跨記憶體位置的順序約束。consumer 看到ready==true時不保證也看得到data=42，會產生資料競爭。producer-consumer 必須用 release store + acquire load（形成 happens-before；RMW 用acq_rel，需全域總順序才用seq_cst）。relaxed 只適合純計數等不需跨位置順序的情境。口訣：scope 最窄、ordering 剛好夠用（別一律 relaxed 也別一律 seq_cst）、位置最近（shared > global）。 - 回讀：Thread Scopes 與 Scoped Atomics
非同步 Barriers 與 Pipelines
Trap: async barrier 的 arrive 與 wait 要分離才有重疊效益
- 陷阱是什麼：把
cuda::barrier當成「比較快的__syncthreads()」，arrive()完馬上wait()，中間不放任何獨立工作。 - 為何容易混淆：async barrier 也是 barrier，直覺以為換上去就自動更快。
- 正確觀念：async barrier 的價值在於把同步切成 arrive 與 wait 兩階段：
bar.arrive()不阻擋 thread、只回傳標記當前 phase 的arrival_token，真正阻擋的是bar.waitmove(token))。效益來自在 arrive 與 wait 之間塞入不依賴他人的獨立工作，把等待延遲藏起來；arrive 即 wait、中間沒工作就退化成普通 barrier，毫無重疊好處（init(&bar, count設定 expected arrival count，arrive 點帶 seq_cst / thread_scope_block 的隱含 fence）。可用性 CC 7.0+，shared-memory 硬體加速與「任意 subset 同步」需 CC 8.0+；注意 cluster remote shared memory 的 barrier 只允許 arrive、不允許 wait，device / system scope 無硬體加速。 - 回讀：非同步 Barriers 與 Pipelines
Trap: 多階段 pipeline vs 單一 barrier 雙緩衝
- 陷阱是什麼：用單一 async barrier 做雙緩衝就以為已達 copy / compute 重疊的極限。
- 為何容易混淆：單 barrier 雙緩衝確實能一邊算一邊載，看起來已經重疊了。
- 正確觀念：單一 barrier 只能做雙緩衝（一個 buffer 在算、一個在載）；
cuda::pipeline多階段 (multi-buffering) 能讓多個 buffer 同時在飛、預取 (prefetch) 後面好幾批，把 copy 延遲更完整地藏在運算之下。pipeline 是 FIFO deque：producerproducer_acquire → producer_commit、consumerconsumer_wait → consumer_release。primitives 版__pipeline_memcpy_async / __pipeline_commit / __pipeline_wait_prior(N)較精簡，但只追蹤 global→shared 複製、有 size/alignment 限制，等價於thread_scope_thread的 pipeline；需要更一般化或不同 thread scope 時改用cuda::pipeline。 - 回讀：非同步 Barriers 與 Pipelines
非同步資料複製與 L1/Shared 配置
Trap:
memcpy_async 需對齊、且搭配 barrier/pipeline 才隱藏延遲
- 陷阱是什麼：呼叫
memcpy_async後沒wait就讀 shared，或以為光是換成memcpy_async就自動隱藏延遲。 - 為何容易混淆：函式名有 async、呼叫也立即返回，看起來資料已就緒。
- 正確觀念：
memcpy_async「彷彿由另一條 thread 執行」global→shared 複製，在 copy 完成前讀寫 shared 或修改 global 都是 data race，必須等wait（或對應 barrier / pipeline 完成）之後才能用。真正隱藏延遲的關鍵是在發起與完成之間做其他有用運算，而非單純改 API（它也順帶繞過同步版shared[i]=global[j]經 register file 中轉的成本）。底層硬體機制 LDGSTS (8.0+, global→shared::cta)、TMA (9.0+, 大型多維 bulk)、STAS (9.0+, register→distributed shared) 各有特定 size / alignment 需求（primitives pipeline 尤其嚴格），不對齊就走不到快路徑。 - 回讀：非同步資料複製與 L1/Shared 配置
Trap: carveout 是 hint（百分比取「下一個更大」）；>48KB shared 需 dynamic opt-in
- 陷阱是什麼：以為
cudaFuncAttributePreferredSharedMemoryCarveout設了百分比就一定照給、設 50% 就剛好拿到 50% shared；又或靜態__shared__陣列直接開超過 48KB。 - 為何容易混淆：carveout 與
cudaFuncSetCacheConfig都在調 L1/shared，容易混為「硬性設定」。 - 正確觀念：L1 與 shared memory 共用 unified data cache。
cudaFuncAttributePreferredSharedMemoryCarveout是 hint，driver 可改用其他配置；當整數百分比對不上支援容量時，會取下一個更大的支援容量（如 CC 12.0、最大 100KB、支援 {0,8,16,32,64,100}KB，設 50% 得 64KB 而非 50KB）。相對地cudaFuncSetCacheConfig是硬性需求，交錯不同 shared 設定的 kernel 會不必要地序列化 launch，故一般偏好cudaFuncSetAttribute。每 block 超過 48KB shared 必須用 dynamic shared memory（extern __shared__），並以cudaFuncAttributeMaxDynamicSharedMemorySize明確 opt-in、在<<<...>>>第三參數傳入大小，且為架構特定。 - 回讀：非同步資料複製與 L1/Shared 配置
CUDA Driver API
Trap: driver API 的 context stack（push / pop 當前 context）
- 陷阱是什麼：以為
cuCtxCreate只是建 context、忘了 host thread 同時只有一個 current context；或函式庫自建 context 後沒還原呼叫者原本的 current context。 - 為何容易混淆：runtime API 隱藏了 context 管理，轉用 driver API 時容易忽略 context 是 per-host-thread 的堆疊狀態。
- 正確觀念：CUDA context 類比一個 CPU process、有獨立 address space（不同 context 的
CUdeviceptr指向不同記憶體）。每個 host thread 維護一個 current context 的堆疊，同時只有堆疊頂為 current：cuCtxCreate/cuCtxPushCurrentpush、cuCtxPopCurrent把 context detach 成 floating 並還原前一個 current。函式庫若自建 context，慣例是用cuCtxPushCurrent/cuCtxPopCurrent包夾、用完還原呼叫者原本的 current context。任何 driver API 呼叫前須先cuInit()；無有效 current context 呼叫相關函式會回CUDA_ERROR_INVALID_CONTEXT。runtime 隱式建立的是 primary context（cuDevicePrimaryCtxRetain()取得），與cuCtxCreate自建 context 不同。 - 回讀：CUDA Driver API
Trap: 要跑未來架構必須載 PTX 而非 binary；buffer 傳參要按 device 端對齊
- 陷阱是什麼：driver API 載入 architecture-specific cubin 卻期待能在更新 GPU 上跑；或用
CU_LAUNCH_PARAM_BUFFER_POINTER傳參時沒按 device 端對齊規則排 offset。 - 為何容易混淆：runtime API 自動處理編譯與傳參，driver API 全要手動，細節容易漏。
- 正確觀念：要在未來架構執行必須載入 PTX（載入時由 driver JIT 成 binary）；cubin（binary）架構特定、不相容未來架構。用單一 buffer 傳 kernel 參數時須符合 device 端對齊：
float4=16、float2=8、CUdeviceptr=__alignof(void*)；device 端double/long long永遠對齊 two-word 邊界（即使 host 用-mno-align-double）；struct 對齊 = 各 field 對齊需求的最大者，device / host 的 padding 可能不同（buffer 以CU_LAUNCH_PARAM_END結尾）。也可改用 pointer 陣列 (args) 傳參較單純。 - 回讀：CUDA Driver API
多 GPU 程式設計
Trap: P2P 需先
cudaDeviceCanAccessPeer 查詢，且 enablePeerAccess 是單向
- 陷阱是什麼：直接呼叫
cudaDeviceEnablePeerAccess而不先查cudaDeviceCanAccessPeer；或在 device 1 啟用對 device 0 的 access 後，就以為 device 0 也能存取 device 1。 - 為何容易混淆：P2P 聽起來像「兩邊互通」，enable 一次就以為雙向都通了。
- 正確觀念：P2P access 取決於 PCIe / NVLink 拓樸，必須先用
cudaDeviceCanAccessPeer()查詢（回 true 才支援）。cudaDeviceEnablePeerAccess()啟用的是「current device → 指定 peer」的單向存取；要雙向需在兩個 device 上各 enable 一次。啟用後同一 UVA 指標可在 peer kernel 直接 deref，但它對該 peer 所有先前與後續分配全域生效，開銷隨 peer 數乘性 (multiplicative) 增長；非 NVSwitch 系統每裝置 peer 連線上限 8 個。更可擴展的做法是用 VMM API 在 allocation time 按需標 peer-accessible。另：peer device memory 的 atomic RMW 僅在單一 GPU 存取該物件時才保證。 - 回讀：多 GPU 程式設計
Trap: multi-GPU current device 語意（stream / event 綁定 device）
- 陷阱是什麼：在 device 1 為 current 時，把綁定 device 0 的 stream
s0拿來 launch kernel；或假設cudaSetDevice之外的呼叫會自動跑在「對的」device。 - 為何容易混淆：stream / event 只是 handle，看不出它綁了哪個 device；current device 是隱含的 host-thread 狀態。
- 正確觀念：current device（
cudaSetDevice設定，首次呼叫前預設 device 0）決定cudaMalloc與 kernel launch 的歸屬；stream 與 event 在建立時即綁定當時的 current device。kernel launch 必須送往綁定 current device 的 stream，否則失敗；但 memory copy 送往非 current device 的 stream 仍成功。cudaEventRecord（event 與 stream 不同 device）與cudaEventElapsedTime（兩 event 不同 device）會失敗；而cudaStreamWaitEvent跨 device 仍成功，可作為跨裝置同步工具。不同 device 的 default stream 之間無順序保證，跨裝置同步屬thread_scope_system。 - 回讀：多 GPU 程式設計
CUDA 功能導覽
Trap: kernel 內 async data copy（shared↔DRAM）別與
cudaMemcpyAsync（CPU↔GPU）混淆
- 陷阱是什麼：看到「asynchronous」就把 kernel 內的 async data copy 與
cudaMemcpyAsync（CPU↔GPU）當同一回事；或不分瓶頸類型亂套功能。 - 為何容易混淆：兩者都叫 async、都在「搬資料」。
- 正確觀念：本章的 async data copies / TMA 指的是 kernel 內部 shared memory ↔ GPU DRAM 的搬移（屬 3.5.1 提升 kernel 效能），不要與 Section 2.5 的
cudaMemcpyAsync（不同元件間 / 相對 CPU 的非同步）混淆。先用功能地圖判斷瓶頸屬於哪一類再選工具：kernel 內效能（async barrier / pipeline / TMA / work stealing）、launch 以上延遲（green context、cudaMallocAsync、CUDA graphs、PDL、lazy loading）、額外能力（EGM、dynamic parallelism）、互通（Direct3D/Vulkan、CUDA IPC）、細緻控制（VMM、driver entry point access、error log）。work stealing 靠 cluster launch control（CC 10.0 Blackwell）；CUDA graphs 可由 stream capture 或 graphs API 建立、適合會重複執行的 workload。 - 回讀：CUDA 功能導覽
第四章：CUDA Features
Purpose
本區延續上方格式，收錄第四章 (CUDA Features) 最容易答錯與最常被混淆的觀念，涵蓋 Unified Memory、CUDA Graphs、Stream-Ordered Allocator、Cooperative Groups、PDL、Green Contexts、Lazy Loading、非同步 barriers/pipelines、三種 async copy (LDGSTS/TMA/STAS)、Work Stealing、L2 Cache Control、Memory Sync Domains、IPC、Virtual Memory Management、EGM、Dynamic Parallelism、Graphics/External Interop 與 Driver Entry Point Access。用法相同：先看標題自我作答，再展開核對「陷阱是什麼 / 為何容易混淆 / 正確觀念」，並沿 回讀 連結深入。
Unified Memory
Trap: hardware-coherent 與 software-coherent 的差異，以及 page size「越大越好」
- 陷阱是什麼：以為 unified memory 的一致性一律靠 page fault + migration、且 virtual page 越大效能越好。
- 為何容易混淆：兩種系統的 programming model「非常相似」，平台差異被抽象掉，看不出底層行為不同。
- 正確觀念：hardware-coherent（如 Grace Hopper）讓 CPU/GPU 共用「邏輯合併的 page table」、一致性粒度為 cache-line、CPU/GPU 並行存取同頁時無需 page fault、競爭更少；software-coherent（Linux HMM 等）CPU/GPU 各自獨立 page table、以 page fault + migration 模擬一致性、粒度為整頁。判別法：合併 page table = hardware、各自 page table = software。Page size 取捨：小頁碎片少但 TLB miss 多、migration 便宜；大頁碎片多但 TLB miss 少、migration 較貴（latency spike 大）；GPU 的 TLB miss 明顯比 CPU 昂貴。調校只針對 virtual page size，別最佳化 physical page size。
- 補充：CUDA IPC 不支援
cudaMallocManaged；但 full-support 系統的 system-allocated memory 具 IPC 能力。對 file-backed 記憶體的 device atomic 只在 hardware-coherent（具hostNativeAtomicSupported）合法，否則 undefined。 - 回讀：Unified Memory：完整支援
Trap:
concurrentManagedAccess = 0 平台 kernel 執行中 CPU 碰 managed data 會 segfault
- 陷阱是什麼：在 Windows / WSL / Tegra（
concurrentManagedAccess = 0）平台上，kernel 還在跑時就讓 CPU 讀寫 managed 變數。 - 為何容易混淆：在 full-support 平台上 CPU/GPU 可並行存取 managed memory，程式碼搬到別的平台「看起來」一樣能跑。
- 正確觀念：這些平台缺 GPU page faulting，任何 kernel 執行期間 GPU 獨佔「所有」managed data（即使該 kernel 根本沒用到那塊），CPU 同時存取 → segmentation fault。正解是先
cudaDeviceSynchronize再存取；或用cudaStreamAttachMemAsync把獨佔粒度從 whole-GPU 縮為 per-stream（未綁定者預設對所有 stream 可見）。這些平台也不可 oversubscribe（上限為 GPU 實體記憶體）。另記：performance hints（cudaMemAdvise/cudaMemPrefetchAsync）永不影響正確性、只影響效能；SetPreferredLocation只是鼓勵、不保證。 - 回讀：Unified Memory：平台與效能提示
CUDA Graphs（擷取與更新）
Trap: stream capture 期間的禁用操作（NULL stream、同步式 API、跨 graph 等待）
- 陷阱是什麼：在 capture 區間呼叫
cudaMemcpy、同步/查詢正被擷取的 stream、用 legacy(NULL) stream、或等待屬於別的 capture graph 的 event。 - 為何容易混淆：capture 把既有 stream 程式碼「錄製」成 graph，讓人以為什麼都能錄。
- 正確觀念：capture 不可用
cudaStreamLegacy(NULL stream)（可用cudaStreamPerThread）；禁止同步或查詢「正被擷取的 stream / captured event / 涵蓋作用中 capture 的 device/context handle」；同步式 API（如cudaMemcpy）會排入 legacy stream 並在返回前同步，故無效；等待屬於「不同 capture graph」的 captured event 無效；等待非擷取 event 需加cudaEventWaitExternal；少數 async 入隊 API（如cudaStreamAttachMemAsync）graph 不支援。CUDA 通則：寧可報錯，也不默默忽略依賴（唯一例外是 stream 進出 capture 模式會切斷前後緊鄰項目的依賴）。違規會 invalidate capture graph，cudaStreamEndCapture仍回傳錯誤值 + NULL graph。 - 回讀：CUDA Graphs：結構與擷取
Trap: graph update 只能改參數，違規就得 fallback 整個重 instantiate
- 陷阱是什麼：以為任何 graph 改動都能用
cudaGraphExecUpdate就地更新。 - 為何容易混淆：update 比 re-instantiate 高效很多，誘使人凡事都想用 update。
- 正確觀念：graph update 只能改「參數」（kernel args、記憶體位址），不能改拓撲、節點類型或「身分」性質的屬性：kernel 不能改 owning context、不能把「不用 dynamic parallelism」改成用；memset/memcpy 只有 1D 可改、不能改 memory type 或
cudaMemcpyKind；external semaphore 不能改數量。拓撲或節點類型一變就必須 re-instantiate（拓撲相關最佳化要重跑）。cudaGraphExecUpdate配對失敗回errorNode與updateResult，只有cudaGraphExecUpdateSuccess才算成功；典型流程是 update 失敗才cudaGraphExecDestroy+cudaGraphInstantiate整個重建。更新於「下一次 launch」生效，不影響進行中的 launch。配對需 API 呼叫順序、dependency 陣列順序、sink node 順序三者一致。 - 回讀：CUDA Graphs：更新與條件節點
CUDA Graphs（條件節點與裝置端啟動）
Trap: conditional node 在 device 評估、與一般 node 不同；SWITCH 越界不是 fallthrough
- 陷阱是什麼：把 conditional node 當普通 node、以為 condition 在 host 評估、或以為 SWITCH 越界會像 C 一樣 fallthrough。
- 為何容易混淆：graph 裡每個 node 在 API 上看起來都差不多。
- 正確觀念：conditional node 的 condition 在 device 上評估（其 dependency 滿足時），把 host CPU 釋放出來。三型：IF（
size=1只有 if body；size=2才有 else，condition==0 執行）、WHILE（進入時評估一次 + 每次 body 完成後再評估）、SWITCH（執行第 n 個 zero-indexed body，condition 不對應任何 body 則「不執行任何 body」，非 fallthrough）。condition value 用cudaGraphConditionalHandle：必須先於節點建立、只能關聯單一 node、無法 destroy；device 端用cudaGraphSetConditional()設值；未指定cudaGraphCondAssignDefault時，每次執行開始 condition 為 undefined。body graph 限制：不可用 Dynamic Parallelism / Device Graph Launch、不可涉及 CUDA arrays、MPS 下不可 cooperative launch。 - 回讀：CUDA Graphs：更新與條件節點
Trap: graph 記憶體節點的固定 VA、銷毀不釋放、與 device graph launch 限制
- 陷阱是什麼：以為含 memory node 的 graph 銷毀會自動釋放記憶體；或以為 device graph 能像 host 一樣重複 launch、能
cudaDeviceSynchronize。 - 為何容易混淆：graph 看似自包含、自帶生命週期管理。
- 正確觀念：graph allocation 的 VA 在整個 graph 生命固定（底層 physical memory 變動不需 graph update），但銷毀含 memory node 的 graph「不會」自動釋放仍存活的配置（須事後
cudaFree/cudaFreeAsync/另一含 free node 的 graph 釋放，否則洩漏；AutoFreeOnLaunch 也不改變此行為）。free 節點必須排在「所有裝置操作完成之後」，kernel 內以記憶體為基礎的同步不足以排序。device graph 須以cudaGraphInstantiateFlagDeviceLaunchinstantiate、只能 host 端 instantiate/update；裝置端對同一 graph 重複 launch →cudaErrorInvalidValue，host+device 同時 launch → undefined；裝置端不能cudaDeviceSynchronize，改用 tail launch（fire-and-forget 上限 120、pending tail launch 上限 255）。 - 回讀：CUDA Graphs：記憶體節點與裝置端啟動
Stream-Ordered Memory Allocator
Trap:
cudaMallocAsync 是 stream-ordered——free 後要「同 stream」才能立即重用、cudaFree 不替你同步
- 陷阱是什麼：以為
cudaFreeAsync後記憶體立刻可被任意 stream 重用、跨 stream 隨意存取/釋放、或用cudaFree釋放 async 配置時 driver 會幫你同步。 - 為何容易混淆：
cudaMallocAsync/cudaFreeAsync命名像普通的 malloc/free，看不出「排入 stream、依 stream 順序生效」的語意。 - 正確觀念：兩者皆 stream-ordered：「配置 → 使用 → 釋放」在同一 stream 是基本正確用法；從「非配置 stream」存取必須保證在配置動作之後，釋放動作開始後再用是 undefined，跨 stream 必須用 event 或 stream 同步接力傳遞「就緒 / 用完」。同一 stream 釋放的記憶體可立即被該 stream 後續配置重用；要供「任意 stream」重用需先讓該 stream 與 CPU 同步。
cudaMallocAsync忽略 current device/context，依「pool 或 supplied stream」決定 device。用cudaFree釋放 async 配置時 driver 不做進一步同步，須自行先cudaStreamSynchronize等避免過早釋放。default(implicit) pool 不支援 IPC；IPC pool 釋放順序 importing 必須先於 exporting。 - 回讀：Stream-Ordered Memory Allocator
Cooperative Groups
Trap: 同步整個 grid 必須用
cudaLaunchCooperativeKernel，一般 <<<>>> 不行
- 陷阱是什麼：寫
this_grid().sync()想同步整個 grid，卻用普通<<<>>>啟動；或把 partition / collective 放進「不是所有 thread 都會到達」的分支。 - 為何容易混淆：
this_grid()隨手可得，看起來就能做全 grid 同步。 - 正確觀念：要同步整個 grid（inter-block 同步）必須用
cudaLaunchCooperativeKernelruntime API 啟動 kernel（保證 launch 為 atomic）；一般<<<>>>只能在 thread block 內同步。支援條件 CC 6.0+（Linux+MPS 需 CC 7.0+），最佳實務先查cudaDevAttrCooperativeLaunch。CUDA 13 起不再支援 multi-device 同步、相關 API 已移除。tiled_partition/labeled_partition/binary_partition與多數 collective 是 collective 操作，全員須參與，放在非全員到達的分支 → deadlock 或 data corruption。memcpy_async只有「來源 global、目的 shared、雙方 ≥4B 對齊」才真的非同步（最佳 16B），且所有 thread 都要先wait才能存取 shared 資料。 - 回讀：Cooperative Groups 深入
Programmatic Dependent Launch
Trap: PDL 的 trigger 與 grid-sync 方向相反、缺一不可；重疊只是「機會」
- 陷阱是什麼：只設了 launch attribute、或只在 secondary 加
cudaGridDependencySynchronize()就以為啟用了重疊；或把程式正確性建立在「primary/secondary 一定並行」上。 - 為何容易混淆：兩個 device 端函式分散在 primary 與 secondary、方向相反，容易漏掉一個。
- 正確觀念：PDL 需三件齊全——primary 端
cudaTriggerProgrammaticLaunchCompletion()（放行訊號）、secondary 端cudaGridDependencySynchronize()（等 primary 結果 flush 到 global memory 的屏障，方向相反、缺一不可）、secondary 用cudaLaunchKernelEx帶cudaLaunchAttributeProgrammaticStreamSerialization。需 CC 9.0+ 才提供真正 overlapping execution。重疊是 opportunistic、不保證並行，依賴並行不安全（可能 deadlock）。若 primary 沒呼叫 trigger，會在「所有 primary thread block 退出後」隱式觸發；不論如何 secondary 仍可能在 primary 資料可見前 launch，故必須做 grid sync。在 Graph 中用cudaGraphDependencyTypeProgrammaticedge：triggerAtBlockStart=0→Programmaticport、=1→LaunchCompletionport。 - 回讀：Programmatic Dependent Launch 深入
Green Contexts
Trap: green context 是「SM 空間切分」不是「時間切分」，且不保證並行
- 陷阱是什麼：把 green context 當成 MPS 那種分時 (time-slicing)，或以為分了 SM 就「保證」兩個 kernel 真的並行。
- 為何容易混淆：GC、MIG、MPS 都是「把 GPU 切給不同工作」，容易混為一談。
- 正確觀念：green context 是空間切分——建立時就綁定固定的、特定的 N 顆 SM（與 work queues），targeting 該 GC 的工作只能用其 provisioned 資源；對比 MIG（靜態切成多個「小 GPU」instance）、MPS（process 級 active-thread 百分比上限，那 N 顆 SM 可任意且隨時間變）。即使分開 provisioned SM 與 WQ「仍不保證」真正並行，只是「移除阻礙並行的因素（降低干擾）」。不需改 kernel，只改 host：建 GC（取資源→split→GenerateDesc→
cudaGreenCtxCreate）+cudaExecutionCtxStreamCreate建屬於 GC 的 stream。CUDA graph 中 node 的 EC 在「建 node 時」決定，graph 被 launch 到的 stream 不決定 SM（僅用於 dependency tracking）。對 primary context 呼叫cudaExecutionCtxSynchronize會連同該 device 上所有 GC 一起同步。 - 回讀：Green Contexts
Lazy Loading 與 Error Log
Trap: lazy loading 預設啟用，但 module 要「首次被用到」才載入
- 陷阱是什麼：以為程式一啟動就把全部 module 載入；或假設 concurrent kernel 一定發生、開機就配滿整片 VRAM、benchmark 不 warmup。
- 為何容易混淆：lazy loading 設計成「無感」(invisible)、預設啟用，從程式碼層面看不出載入時機被延後。
- 正確觀念：CUDA 12.3 起所有平台預設啟用 lazy loading（
CUDA_MODULE_LOADING=LAZY），module 等到真正被需要時才載入；需 runtime ≥ 11.7 AND driver ≥ 515（缺一即全部回退 eager），不需 compiler 支援（pre-11.7 的 SASS/PTX 也能 lazy load）；含 managed variable 的 module 仍 eager 載入。三大潛在陷阱：(1) 程式假設 concurrent kernel 但執行被序列化 → deadlock；(2) 開機吃滿 VRAM 使延後載入的 module 配不到記憶體；(3) 初始化被算進被量測視窗污染 benchmark。萬用解法是「preload 你要用的 kernel」（cuModuleGetFunction()/cudaFuncGetAttributes()會觸發載入，cuModuleLoad()不保證立即載入），或CUDA_MODULE_LOADING=EAGER、用cudaMallocAsync、做 warmup 迭代。查 mode 用cuModuleGetLoadingMode（須先cuInit）。 - 回讀：Lazy Loading 與 Error Log
Asynchronous Barriers 與 Pipelines
Trap: async barrier 要把 arrive 與 wait「分離」並塞入獨立工作才有重疊效益
- 陷阱是什麼：把
cuda::barrier當成「比較快的__syncthreads()」，arrive()完馬上wait()，中間不放任何獨立工作。 - 為何容易混淆：它也是 barrier，直覺以為換上去就自動更快。
- 正確觀念：價值在把同步切成 arrive 與 wait 兩階段：
bar.arrive()不阻塞 thread、只回傳綁定「當前 phase」的arrival_token，真正阻塞的是bar.waitmove(token))；效益來自在 arrive 與 wait 之間塞入不依賴他人的獨立工作隱藏延遲，arrive 即 wait、中間沒工作就退化成普通 barrier。init(&bar, count)設 expected arrival count，必須在任何 threadarrive()之前用block.sync(bootstrap。token 只能用「當前或緊接前一個 phase」否則 undefined；wait只保證「全員已 arrive」、不保證全員已 wait；arrive 與 wait 之間不可用任何 collective。warp 完全收斂只更新 barrier 1 次、完全發散 32 次（建議__syncwarp收斂後再 arrive）。transaction barrier（CC 9.0+、限 shared 的 block/cluster scope）的 wait 需同時滿足「全員 arrive」且「transaction count 達標」。 - 回讀：Asynchronous Barriers 深入
Trap: pipeline 的
count 是「緩衝深度」非總 batch 數；partitioned 與 warp diverge 的隱藏成本
- 陷阱是什麼：把
pipeline_shared_state<scope, count>的count當成「總 batch 數」；以為 unified pipeline 沒成本；diverge 後不收斂照樣 commit。 - 為何容易混淆：pipeline 高階 API 把底層 barrier 與糾纏行為都藏起來了。
- 正確觀念：
count是「能同時在飛的 stage 數（緩衝深度）」，不是總 batch 數；非thread_scope_thread的 scope 才需要pipeline_shared_state。partitioned（固定 producer/consumer 角色、不可變）pipeline 每個 stage 配一組 shared memory barriers，unified 也付這些成本（其實可改用__syncthreads()），能用 thread-local pipeline 就優先用。producer_acquire()在資源用盡時阻塞（形成背壓）。warp diverge 時 thread 的「感知序列」< 「實際序列」會 over-wait（極端 diverge 每 thread 可能等 32 個 batch），故 commit 前先__syncwarp收斂。提前離開的 thread 必須先cuda::pipeline::quit()，否則其餘 thread 在集體操作上永久等待。 - 回讀：Pipelines 深入
非同步資料複製（LDGSTS / TMA / STAS）
Trap: LDGSTS vs TMA vs STAS 的方向、CC 門檻與 async-group 不可混用
- 陷阱是什麼：以為三種 async copy 方向可互換（例如用 LDGSTS 做 shared→global、用 STAS 碰 global），或混用兩套 wait_group。
- 為何容易混淆：三者都叫 async copy、都用 mbarrier 標完成。
- 正確觀念：方向與用途各不同：LDGSTS（CC 8.0+）只 global → shared::cta，適合小、element-wise 傳輸、繞過 register（4/8B 走 L1 ACCESS、16B 走 L1 BYPASS，最佳 128B 對齊，預設每 thread 只等自己的複製）；TMA（CC 9.0+）做 global ↔ shared 的大型/多維 bulk（讀＝global→shared 用 shared memory barrier 的 transaction count、block 內任一 thread 可等；寫＝shared→global 用 bulk async-group
commit/wait、僅發起 thread 可等）；STAS（CC 9.0+）是唯一 register → distributed shared memory（cluster 內跨 block，一次 4/8/16B，只有低階cuda::ptx::st_async、無高階包裝）。bulk 與 non-bulk 的 async-group 獨立：cp.async.wait_group(LDGSTS) 與cp.async.bulk.wait_group(TMA) 不可混用。TMA 應由單一 thread 用is_elected()/invoke_one發起，別只寫if(threadIdx.x==0)（編譯器可能插入 peeling loop 造成 warp serialization）。 - 回讀：非同步複製：LDGSTS、非同步複製：STAS
Trap: 多維 TMA「必須 host 端 encode tensor map」，且 global 的 tensor map 需 release-acquire
- 陷阱是什麼：以為多維 TMA 能像 1D bulk 一樣只給指標與大小就搬；或在 device 隨意改 tensor map、或用 global memory 的 tensor map 卻不做 proxy fence。
- 為何容易混淆：1D bulk-asynchronous copy 不需 tensor map，容易推論「多維也不用」。
- 正確觀念：多維（最多 5D）TMA 必須在 host 端用
cuTensorMapEncodeTiled（driver API）建立CUtensorMap描述陣列 layout，device 端只給 tile 左上角座標{x, y, …}，硬體自動算 global 位址、讀越界自動 zero-fill（座標可為負；寫可部分越界但座標不可負）。約束：最快變動維排第一、stride 須 16B 倍數、多維 shared 須 128B 對齊（swizzle 128B 時 1024B）。傳遞首選const __grid_constant__參數，其次__constant__，再者 global memory（每 block 用前需fence_proxy_tensormap_generic）。device 端編碼僅支援 tiled-type、標記sm_90a專屬、.rank欄位零基（填 desired_rank − 1）。用 global memory 的 tensor map 必須在「修改方 release」（tensormap_cp_fenceproxy）與「使用方 acquire」之間建立 proxy 順序，且 fence 與 use 須在同一 block。計算經 generic proxy 寫 shared、TMA 經 async proxy 讀，中間需fence_proxy_async。 - 回讀：非同步複製：TMA
Work Stealing 與 Cluster Launch Control
Trap: CLC 取消失敗後「觀察再 request」與「取用失敗 index」都是 UB
- 陷阱是什麼：以為 try-cancel 失敗後還能再發 request、或去讀失敗 request 的 block index。
- 為何容易混淆：
try_cancel看起來像普通的 try-then-retry。 - 正確觀念：Cluster Launch Control（Blackwell, CC 10.0）讓 block 取消「尚未開始執行」的 block，成功則竊取其 index 做事，即 work stealing，結合 Fixed Work（load balancing / preemption）與 Fixed Number（reduced overhead）的優點。兩條 UB 規則：(1) 觀察到失敗（即 query 結果）之後再發 request 是 UB——關鍵在「兩個 request 之間有沒有插入對結果的 query 觀察」，而非是否失敗；(2) 取用失敗 request 的 block index 是 UB（只有 success 時 index 才有效）。CLC 是 async proxy 操作，需
fence_proxy_async_generic_sync_restrict的 acquire/release；try_cancel是 uniform 指令，建議包進invoke_one。cluster 版用try_cancel_multicast、scope 全提升為scope_cluster、取消前所有 block 須存在（cluster_group::sync()），解出的bx還要加block_index().x。 - 回讀：Work Stealing 與 Cluster Launch Control
L2 Cache Control
Trap: L2 set-aside 開太大、
hitRatio=1.0 反而傷一般存取；用完要主動 reset
- 陷阱是什麼：以為 persisting set-aside 撥越大越好、
hitRatio一律設 1.0、且用完不 reset。 - 為何容易混淆：「保留更多 cache 給熱資料」直覺上總是有益。
- 正確觀念：set-aside 從 L2 撥給 persisting 存取優先使用，但保留區被「所有並行 kernel 共用」、總使用量是各 kernel 加總，超過容量時持久化好處遞減，且過大的 set-aside 會壓縮 normal/streaming 可用的 L2、傷一般存取。
hitRatio < 1.0可避免 thrashing：set-aside 16KB、window 32KB 時，hitRatio=1.0會嘗試把整個 32KB 都快取、不斷互相驅逐 (thrash)，0.5只標 16KB 為 persisting、剛好放下不抖動。用完 persisting 務必先把 windownum_bytes設 0 關窗、再呼叫cudaCtxResetPersistingL2Cache()，否則 persisting line 長期占住 L2（不要依賴自動 reset，時間不確定）。需 CC 8.0+；MIG 模式下停用；MPS 下不能用cudaDeviceSetLimit改大小（須 MPS server 啟動時用環境變數CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT）。 - 回讀：L2 Cache Control
Memory Synchronization Domains
Trap: memory sync domain 是為了「隔離 fence 干擾」，跨 domain 需 system-scope fence
- 陷阱是什麼：不知道 fence 會「等待超出必要的 in-flight transactions」而拖慢 compute kernel；或以為把通訊 kernel 換到別的 domain 後，跨 domain 仍能用 device-scope fence 同步。
- 為何容易混淆：fence interference 是隱性的效能退化，從正確性角度看不出問題。
- 正確觀念：memory fence 因 cumulativity（累積性）——GPU 無法分辨某個 in-flight write 是 source level 真正保證要可見的、還是只是碰巧時序可見，只能保守地等待全部——而 over-wait；典型是 local compute kernel 完成時隱式 flush，被並行的 NCCL 慢的 NVLink/PCIe 寫入拖累。Memory Synchronization Domains（CC 9.0 Hopper + CUDA 12.0，Hopper 有 4 個 physical domain）讓每個 launch 取得 domain ID，write 與 fence 都標記該 ID、fence 只 order 同 domain 的 write；把通訊 kernel 放到不同 domain 即可隔離流量。跨 domain 的 ordering 必須用 system-scope fence，同 domain 內 device-scope 即足夠。用兩個 launch attribute：
MemSyncDomain選 logical（Default/Remote）、MemSyncDomainMap設 logical→physical 映射；預設 default→0、remote→1，kernel 預設落 domain 0 維持向後相容（pre-9.0 回報 count=1）；NCCL 2.16 起自動標 remote。選 domain 不改變 kernel 可做的記憶體存取、只影響 fence ordering。 - 回讀：Memory Synchronization Domains
Interprocess Communication
Trap: 跨 process 要交換「handle」不是 pointer；
cudaMalloc 子分配會連帶洩漏整塊
- 陷阱是什麼：把原始 device pointer 直接傳給另一個 process；或用
cudaIpcGetMemHandle分享cudaMalloc的小配置，不知會把整個底層區塊都分享出去。 - 為何容易混淆：pointer 看起來就是個「位址」，直覺以為跨 process 也能用。
- 正確觀念：device pointer 與 event handle 只在「建立它的 process」內有效；跨 process 必須交換 handle——
cudaIpcGetMemHandle()取得 → 經標準 OS IPC（shared memory / 檔案）傳遞 →cudaIpcOpenMemHandle()還原成對方的 process-local pointer。永遠交換 handle、由對方重新 open，切勿直接傳 pointer。Legacy IPC API 僅 Linux、不支援cudaMallocManaged、收發雙方 driver/runtime 須一致；cudaMalloc可能從較大區塊子分配，IPC 會分享「整個底層 block」可能洩漏其他子分配，建議只分享 2 MiB 對齊的配置；Tegra/L4T 只支援 event-sharing。VMM IPC 可逐 allocation 控制 peer 存取與分享、跨多種 OS（須改用 Driver API）。多節點 NVLink 叢集用 fabric handle。 - 回讀：Interprocess Communication
Virtual Memory Management
Trap: VMM 的
reserve → map → setAccess 順序，與 unmap → release → addressFree 釋放順序
- 陷阱是什麼：以為
cuMemCreate回傳的 handle 就是可用指標、cuMemMap之後就能直接被 kernel 存取；或釋放時把順序顛倒。 - 為何容易混淆：習慣了
cudaMalloc一步到位回傳可用指標。 - 正確觀念：VMM 把「虛擬位址保留」與「實體配置」分離，須顯式三步：
cuMemCreate配實體（回傳CUmemGenericAllocationHandle，「非指標、尚不可存取」，size 須對齊cuMemGetAllocationGranularity）→cuMemAddressReserve保留 VA（概念等同mmap/VirtualAlloc）→cuMemMap把 VA 綁到實體 handle →cuMemSetAccess必做：只 map 不設 access，被 kernel 存取會 crash。釋放嚴格依序cuMemUnmap → cuMemRelease → cuMemAddressFree（顛倒會出錯；OS-specific handle 還需fclose、fabric 不需）。動機：傳統cudaEnablePeerAccess會把「所有過去與未來」配置都映射到 peer，VMM 改為 allocation 粒度控制。整套屬 Driver API、需系統支援 UVA；fabric handle（CU_MEM_HANDLE_TYPE_FABRIC）跨節點需 IMEX。multicast 須所有 device 先cuMulticastAddDevice才能cuMulticastBindMem。 - 回讀：Virtual Memory Management
Extended GPU Memory
Trap: EGM 用 NUMA node id 而非 device ordinal；別用 cgroups 限制裝置
- 陷阱是什麼：用 device ordinal 當 EGM 配置的 location id；或用 cgroups 限制可見裝置。
- 為何容易混淆：平常定位裝置都用 device ordinal。
- 正確觀念：EGM 用 OS 指派的 NUMA node id (numaID) 表達記憶體放置位置，它「不等於」device ordinal（關聯到「最近的 host node」），用
cuDeviceGetAttribute+CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID取得。支援兩種 allocator：cuMemCreate(VMM) 與cudaMemPoolCreate(Stream Ordered)，location type 用CU_MEM_LOCATION_TYPE_HOST_NUMA/cudaMemLocationTypeHostNuma。多節點需CU_MEM_HANDLE_TYPE_FABRIC+ export/import shareable handle（經 TCP/IP），映射後仍須對每顆本地 GPU 逐一cuMemSetAccess。別用 cgroups 限制裝置（會阻斷 EGM 路由、掉效能），改用CUDA_VISIBLE_DEVICES。EGM 以 2MB pages 映射，超大配置可能更多 TLB miss；本機存取走 NVLink-C2C，把遠端 socket 記憶體映射為 EGM 反而更快（流量保證走 NVLink）。 - 回讀：Extended GPU Memory
CUDA Dynamic Parallelism
Trap: CDP2 沒有
cudaDeviceSynchronize，不能從 parent/child 顯式同步 child
- 陷阱是什麼：在 device code 用
cudaDeviceSynchronize()等 child 完成、或想在 child 內再同步它啟動的 child；把 local/shared 指標傳給 child。 - 為何容易混淆：沿用 CDP1 與 host runtime 的同步習慣。
- 正確觀念：CDP2（CUDA 12.0+ 預設、CC 9.0+ 唯一）移除了
cudaDeviceSynchronize()（device code 引用會編譯錯誤）——無法從 parent thread 顯式同步 child，因此不保證 child 的修改對 parent 可見；parent 退出前取回 child 結果的唯一方式是把 kernel 送進cudaStreamTailLaunchstream（另有cudaStreamFireAndForget）。global memory 是 weak consistency，唯一完全一致點 = child 被 invoke 的那一刻（launch 前 parent 的寫入對 child 可見）。傳給 child 的指標必須是 global/mapped（用__isGlobal()判斷）；把 local/shared 指標傳給 child，或傳給cudaMemcpy*Async/cudaMemset*Async是非法的。CDP 不引入任何新的並行保證（block 間、parent/child 間都不可依賴並行）。device 端建立的 stream/event 只在「建立它的 grid」內有效。編譯需-rdc=true -lcudadevrt；CDP1 與 CDP2 互相 launch / 同一 call graph 混用 →cudaErrorCdpVersionMismatch。 - 回讀：CUDA Dynamic Parallelism
Graphics / External Resource Interop
Trap: 匯入外部資源需「相符的裝置」(UUID/LUID)、destroy 不釋放映射、handle 所有權因平台而異
- 陷阱是什麼：在「不同裝置」匯入外部資源；以為
cudaDestroyExternalMemory會把映射也釋放；Linux fd 匯入後還繼續使用它。 - 為何容易混淆：import 看起來只要 handle 對就能用，所有權細節容易被忽略。
- 正確觀念：匯出的資源必須在「建立它的同一裝置」匯入：Vulkan 比對 device UUID、Direct3D12 比對 device LUID、NVSCI 用 GPU id（
CUuuid）；且 Vulkan physical device group count 須為 1、D3D12GetNodeCount須為 1。cudaImportExternalMemory後須cudaExternalMemoryGetMappedBuffer/MipmappedArray映射，offset/size 等必須與匯出端設定一致否則 UB；cudaDestroyExternalMemory不釋放映射（device pointer 另以cudaFree、array 另以cudaFreeMipmappedArray釋放）。handle 所有權：Linux fd 匯入後 CUDA 接管（勿再用）、Windows NT handle 由應用程式自行CloseHandle、D3DKMT 自動銷毀；匯入 D3D12 資源一律須設cudaExternalMemoryDedicated。semaphore：binary 的 wait 必須在對應 signal「已發出之後」才能發；D3D12 fence 單調遞增、wait 等到value ≥ 指定值。另：直接 Graphics Interop 只支援 D3D9/10/11、不含 D3D12，須先 register 再取 bindless handle，map 期間視為 CUDA 獨佔（其間用 GL/D3D/另一 context 存取為 undefined）。 - 回讀：External Resource Interop、Graphics Interoperability
Driver Entry Point Access
Trap:
cuGetProcAddress 的 version 要「精確對應 typedef」，且兩種失敗碼意義不同
- 陷阱是什麼：把
CUDA_VERSION或cuDriverGetVersion()回傳值當cuGetProcAddress的 version 引數；或為「方便」傳更高版本；分不清「版本不夠」與「找不到符號」。 - 為何容易混淆：version 引數看起來就是「我現在用的 CUDA 版本」。
- 正確觀念：
cuGetProcAddress(name, &pfn, version, flags, &driverStatus)的 version 必須「精確對應你所用 typedef 的版本號」（PFN_xxx_vNNNNN的 NNNNN = 引入該符號的 CUDA 版本，如 3.2=3020、10.1=10010），應硬編常數、別用CUDA_VERSION或cuDriverGetVersion；傳「更高」版本未來可能換回更新的符號 → ABI/簽章不符的 undefined behavior；版本低於符號引入版本 →CUDA_ERROR_NOT_FOUND。呼叫前先用cuDriverGetVersion確認 driver 夠新。兩類失敗分管道：(1) API/usage 錯誤（pfn=NULL、非法 flags）→CUresult回傳碼；(2) 找不到符號 →driverStatus（CUdriverProcAddressQueryResult）：VERSION_NOT_SUFFICIENT= 「你的 cudaVersion 引數太低、driver 其實夠新（升 cudaVersion 即可）」、SYMBOL_NOT_FOUND= 「driver 太舊或名稱拼錯，cudaVersion 給多少都無關」。自 CUDA 11.3 起提供，類比dlsym/GetProcAddress；per-thread default stream 版本符號帶_ptsz/_ptds（影響同步行為），可用取址 flags 強制。 - 回讀：Driver Entry Point Access