面試陷阱題 (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
第五章：Technical Appendices
Purpose
本區延續上方格式，收錄第五章 (Technical Appendices) 最容易答錯與最常被混淆的觀念，涵蓋 compute capability 與環境變數、device code 的 C++ 標準支援與語言限制、C 標準函式庫與 lambda、函式/變數限定詞與啟動界限、同步原語/原子/記憶體屏障、warp 函式與巨集、編譯器提示與 Warp Matrix (WMMA)、浮點與數學函式（standard vs intrinsic）、裝置端 runtime (CDP2)，以及 CUDA C++ 記憶體模型與執行模型。用法相同：先看標題自我作答，再展開核對「陷阱是什麼 / 為何容易混淆 / 正確觀念」，並沿 回讀 連結深入。
計算能力與環境變數
Trap: compute_NNa / compute_NNf / compute_NN 的執行相容性完全不同
- 陷阱是什麼：以為帶
a（architecture-specific）或f（family-specific）後綴編出來的程式，能像一般 baseline 那樣「在更高 CC 上也能跑」。 - 為何容易混淆：三者只差一個字母，又同屬一組 feature set，直覺以為相容範圍一致。
- 正確觀念：三組 feature set 為含括關係 architecture-specific ⊇ family-specific ⊇ baseline。
compute_NN（baseline，無 arch 專屬特性）相容 CC NN.x 及之後所有架構；compute_NNf（CC 10.0 起、後綴f）只在同 family 裝置可執行（如compute_100f相容 10.0 與 10.3、compute_120f相容 12.0 與 12.1）；compute_NNa（CC 9.0 起、後綴a）只能在編譯時那個 exact CC 上執行。許多上限跨所有 CC 不變：warp size=32、max threads/block=1024、max 32-bit registers/thread=255。執行期查 CC 用cudaDeviceGetAttribute(cudaDevAttrComputeCapabilityMajor/Minor) /cuDeviceGetAttribute/nvmlDeviceGetCudaComputeCapability（NVML 需-lnvidia-ml），命令列用nvidia-smi --query-gpu=name,compute_cap。 - 回讀：計算能力
Trap: CUDA_VISIBLE_DEVICES 遇無效 index 會「截斷」清單，且會重排 ordinal
- 陷阱是什麼：以為
CUDA_VISIBLE_DEVICES只是過濾清單、順序無關緊要；或以為清單裡夾一個無效值只會略過該值。 - 為何容易混淆：把它當成單純的「白名單」，沒注意到它同時決定列舉順序與 ordinal。
- 正確觀念：未設定＝所有 GPU 可見；設為空字串＝沒有任何 GPU 可見。取值可用 integer index、GPU UUID（可用唯一前綴縮寫）、MIG 識別（僅支援單一 MIG instance）。遇到無效 index 會截斷：
CUDA_VISIBLE_DEVICES=0,2,-1,1只有 device 0 與 2 可見（排在無效-1之後的都不可見）。列舉順序決定 ordinal：=2,1時cudaSetDevice(0)設的是實體 device 2（它列舉在最前、ordinal 為 0），且cudaGetDeviceCount()只計入可見 device，整數 device id 只接受[0, 可見數-1]。列舉順序另受CUDA_DEVICE_ORDER影響（預設FASTEST_FIRST，另有PCI_BUS_ID）。 - 回讀：CUDA 環境變數
Trap: CUDA_LAUNCH_BLOCKING 是除錯工具、CUDA_MODULE_LOADING 預設 LAZY
- 陷阱是什麼：把
CUDA_LAUNCH_BLOCKING=1當成效能設定常開；或以為 module（kernel）預設在程式啟動時就全部載入。 - 為何容易混淆：「blocking」聽起來像同步化能避免競爭；而靜態連結的印象讓人以為 kernel 一開始就全載好。
- 正確觀念：
CUDA_LAUNCH_BLOCKING=1關閉非同步執行、讓 GPU 工作從 CPU 視角同步，價值是「讓 CUDA API 錯誤剛好在觸發它的那個呼叫被觀察到」以利除錯，會讓程式變慢，平時不應開啟。CUDA_MODULE_LOADING預設為 LAZY：延後到取得 function handle 才載入，降低 startup 時間與 GPU 記憶體佔用、首次呼叫才載入；EAGER則初始化即全載，startup 與記憶體較高、但 launch 開銷可預測。另記 JIT 快取：CUDA_CACHE_DISABLE=1每次重編 PTX→CUBIN；預設大小桌機/伺服器 1 GiB、嵌入式 256 MiB、上限 4 GiB；CUDA_FORCE_PTX_JIT覆蓋CUDA_FORCE_JIT、CUDA_DISABLE_PTX_JIT覆蓋CUDA_DISABLE_JIT。 - 回讀：CUDA 環境變數
C++ 語言支援與限制
Trap: device code 不是完整 ISO C++——RTTI、例外、concurrency、Modules/Coroutines 都不支援
- 陷阱是什麼：在 kernel /
__device__函式裡用try/catch/throw、typeid/dynamic_cast、std::atomic/thread_local、long double，或期待 C++20 的 Modules、Coroutines 可用。 - 為何容易混淆：nvcc 宣稱支援 C++03/11/14/17/20/23 全部標準的「所有語言特性」，讓人以為 device code 也是完整 ISO C++。
- 正確觀念：device code 不支援 RTTI（
typeid、dynamic_cast）、例外（try/catch/throw）、long double、trigraphs。C++11 整個 Concurrency 群組（std::atomic、memory model、bidirectional fences、thread_local、propagating exceptions）在 device 都不支援，CUDA 改用自家原語與 libcu++ 的cuda::atomic。各標準起始 Toolkit：C++11=7.x、C++14=9.x、C++17=11.x、C++20=12.x。C++14 只缺 N3664/N3778（兩項 operator new/delete 配置）；C++17 列表內語言特性全支援（structured bindings、constexpr if、constexpr lambda、fold expressions、inline variables），是現代 CUDA 的安全基準；C++20 的 Concepts、operator<=>、consteval、constinit、char8_t皆可用，唯 Modules 與 Coroutines 不支援。需要 STL 設施時用 host/device 皆可用的 libcu++（cuda::std::）。C++23 device code 多數可用（Deducing this、if consteval、[[assume]]），但有 5 項不支援（如 narrowing-contextual-conversions-to-bool、部分 constexpr 放寬）；host compiler 版本門檻：C++20 需 GCC ≥ 10 / Clang ≥ 10 / MSVS ≥ 2022、C++23 需 GCC ≥ 14 / Clang ≥ 18，而 MSVC 對 C++23 為 Not Supported。 - 回讀：C++ 語言特性支援
Trap:
__global__ 參數有 32,764 bytes 上限、用 raw memcpy 傳遞、跨空間解指標是 UB
- 陷阱是什麼：傳一個帶 user-defined copy constructor / non-trivial destructor 的大物件給 kernel 並指望語意如標準 C++；或在 host 端解參考 device 指標（反之亦然）。
- 為何容易混淆：kernel 呼叫長得像普通函式呼叫，讓人以為參數傳遞、複製語意、指標解參考都照 ISO C++。
- 正確觀念：
__global__參數經 constant memory 傳遞，總大小上限 32,764 bytes，且不可 variadic（C.../va_list）、不可為std::initializer_list、polymorphic class 參數為 UB。從 host 啟動時，CUDA Runtime 用 rawmemcpy複製 raw 記憶體，user-defined copy constructor 被跳過（其副作用不發生）；又因 launch 非同步，non-trivial destructor 可能在 kernel 完成前就在 host 執行。從 device 啟動則每個 argument 必須 trivially copyable 且 trivially destructible。指標解參考只允許在記憶體所屬 execution space 內，host 解 device 指標、device 解 host 指標皆 UB（多半 segfault）；cudaGetSymbolAddress()取得的位址只能用於 host code。__global__/__tile_global__/__tile__不支援 recursion，且__global__不可為 class member、不可宣告為constexpr、不可有auto推導回傳型別。 - 回讀：C/C++ 語言限制
C 標準函式庫、Lambda 與多型包裝器
Trap: 標了
__device__ 不一定是 extended lambda；*this 捕捉與 nvstd::function 邊界
- 陷阱是什麼：以為只要在 lambda 上標
__device__就是 extended lambda 並能傳進 kernel；或在 host 函式定義的 device lambda 直接存取 class member（this->member）；或把nvstd::function從 host 傳到 device。 - 為何容易混淆：執行空間標註看起來就決定一切，且 lambda 捕捉 member 與一般 C++ 寫法無異。
- 正確觀念：lambda 執行空間由最內層 enclosing function 的執行空間決定（無 scope 則為
__host__）；要傳進__global__必須是__device__或__host__ __device__，且 global/namespace scope 的 lambda 不可當引數。extended lambda 必須定義在__host__或__host__ __device__函式內並帶執行空間標註（需--extended-lambda，定義__CUDACC_EXTENDED_LAMBDA__）；因此__device__/__global__函式內或 global scope 的標註 lambda 都不是 extended lambda。捕捉限制：只能 by value、陣列維度 ≤ 7、host-device 版不可 init-capture、不可捕 variadic pack；且 extended lambda 的存在/順序不可依賴__CUDA_ARCH__。C++11/14 規則下 lambda 捕的是this指標，於 GPU 存取 hostthis->member會 runtime error；C++17*this捕捉複製物件本身來解決。nvstd::function（<nvfunctional>）host/device 皆可用但不可跨 host/device 邊界傳遞或初始化，其operator()僅__device__。 - 回讀：C 標準函式庫與 Lambda
Trap: device printf 回傳引數數、buffer 程式結束不 flush；device malloc 不可用 cudaFree
- 陷阱是什麼：以為 device
printf()像標準 C 回傳字元數、輸出一定會出現；或以為 device 端malloc()配的記憶體能用cudaFree/cudaMemcpy操作。 - 為何容易混淆：函式名與標準 C 一模一樣，heap 配置也都叫
malloc，自然套用主機端直覺。 - 正確觀念：device
printf()每個 thread 各自執行，回傳解析的引數數（非字元數）：無引數→0、format為NULL→-1、內部錯誤→-2；引數上限 32 個（多出的被忽略）；輸出在 host 端格式化。輸出寫進固定大小的環狀 buffer，預設 1 MB（cudaLimitPrintfFifoSize可調），只在 kernel launch、同步、blocking memcpy、module load/unload、context 銷毀、host callback 前 flush，程式結束不會自動 flush。devicemalloc()/new從 device heap（預設 8 MB，須在任何配置前用cudaDeviceSetLimit(cudaLimitMallocHeapSize, …)設定、module 載入後不可變）配置，回傳 16-byte 對齊；device heap 與 host runtime API 記憶體互不相通——device 配的不可用cudaMalloc/cudaMemcpy/cudaMemset/cudaFree操作，反之亦然。 - 回讀：C 標準函式庫與 Lambda
函式與變數限定詞
Trap:
__launch_bounds__ 的兩個參數與 register 取捨方向相反
- 陷阱是什麼：以為
__launch_bounds__永遠是「壓低暫存器」；或忽略只給maxThreadsPerBlockvs 同時給minBlocksPerMultiprocessor的差別；或把它和__maxnreg__一起套。 - 為何容易混淆：兩個數字都和「占用率/暫存器」有關，方向卻不直覺。
- 正確觀念：
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor[, maxBlocksPerCluster])，maxThreadsPerBlock必要、編成 PTX.maxntid、minBlocksPerMultiprocessor編成.minnctapersm。編譯器據此推導暫存器上限 L 以確保minBlocksPerMultiprocessor個 block 能常駐；初始用量超過 L 就壓低（可能多 local memory / 指令）。但若初始用量低於 L：只給maxThreadsPerBlock時用它決定 n↔n+1 常駐 block 的門檻；同時給minBlocksPerMultiprocessor時編譯器可能反而把用量提高到 L，以減少指令數、改善延遲隱藏。以超過maxThreadsPerBlock的 threads/block 啟動會失敗（「too many resources requested for launch」），故建議至少加單引數版。__launch_bounds__與__maxnreg__（編成.maxnreg）不可同時套；含__maxnreg__的 kernel 會忽略--maxrregcount。最佳值常隨架構不同，可用__CUDA_ARCH__在 device 端切換（host 端__CUDA_ARCH__未定義）。 - 回讀：函式與變數限定詞
Trap:
__restrict__ 會增暫存器壓力、__grid_constant__ 唯讀、int3 對齊不是 16
- 陷阱是什麼：以為
__restrict__只有好處沒有代價；以為__grid_constant__參數可改；或假設int3/float3像int4一樣是 16-byte 對齊。 - 為何容易混淆：
__restrict__直覺是「放開優化＝一定更快」，向量型別則容易以為對齊都等於 size 補齊到 16。 - 正確觀念：
__restrict__承諾指標不別名（no aliasing），所有指標參數都要加才有效；副作用是把載入與 CSE 快取到暫存器、增加暫存器壓力，可能降 occupancy 反傷效能；標了__restrict__的__global__const 指標會編成唯讀快取載入（PTXld.global.nc，等同__ldg()）。__grid_constant__讓__global__參數全 thread 共用單一位址、唯讀（修改含 mutable 成員為 UB）、具 kernel 生命週期，參數須為 const-qualified 非 reference。向量型別對齊：int4/uint4/float4/double2/longlong2為 size 16、align 16，但int3/float3為 size 12、align 4（非 16）。gridDim/blockDim是dim3、blockIdx/threadIdx是uint3、warpSize是執行期int（常見 32）；__cluster_dims__需 CC 9.0+。CUDA 13 已 deprecatedlong4/ulong4/longlong4/ulonglong4/double4。 - 回讀：函式與變數限定詞
Trap: 記憶體空間限定詞的存取者/生命週期，與
__managed__ decltype、static __shared__ 不可初始化
- 陷阱是什麼：搞混
__device__/__constant__/__shared__/__managed__的可存取者與生命週期；以為__managed__變數位址是常數運算式；在 static__shared__宣告處給初值；或對 device 全域符號直接用&/memcpy跨 host/device 搬資料。 - 為何容易混淆：四個限定詞都「放在 device 上」、又都是全域變數語法，容易以為存取規則一致。
- 正確觀念：
__device__變數在 device 全域記憶體、grid 內 threads 與 Runtime API 可存取、生命週期為 program/CUDA context、每 device 一份；__constant__同位置但 device 端唯讀、只能由 host 改；__shared__在 SM 上、僅同 block threads 可存取、生命週期僅該 block、每 block 一份；__managed__由 host/device threads 自動共用、生命週期為 program、每 program 一份。host 端用cudaGetSymbolAddress/cudaGetSymbolSize/cudaMemcpyToSymbol/cudaMemcpyFromSymbol存取__device__/__constant__符號（取得的位址只能用於 host code），不可直接對 device symbol 解參考或memcpy。__managed__三大限制：位址不是常數運算式（不能當 template 非型別參數）、不可為 reference 型別、runtime 未就緒時（static/dynamic init/destruct、exit()後）不可用其位址或值；且decltype(global_var)是錯誤（非括號 id-expression），decltype((global_var))才合法。__shared__尺寸分 static（編譯期）與 dynamic（extern __shared__陣列/指標，大小由<<<,,bytes>>>給），static 尺寸變數不可在宣告處初始化。 - 回讀：函式與變數限定詞
同步、原子與記憶體屏障
Trap:
__threadfence 與 __syncthreads 不同——fence 只排序、不集合、不保證可見性
- 陷阱是什麼：把
__threadfence()當成「跨 block 的 barrier」來等其他 block；或以為 fence 過後寫入就一定對別人可見。 - 為何容易混淆：兩者名字都像「同步」，又都和記憶體一致性有關。
- 正確觀念：
__syncthreads*()是 block 層級 barrier，等 block 內所有未退出 thread 到達同一呼叫點，並建立 strongly happens-before 的記憶體排序；放在條件式中時條件須對整個 block uniform，否則可能 hang／未定義（if(threadIdx.x>0) __syncthreads()是錯的）；變體_count/_and/_or在同步之餘對 predicate 做規約。memory fence 不集合、不等待、只排序：__threadfence_block（block scope）、__threadfence（device scope）、__threadfence_system（system：device 內所有 thread + host thread + 所有 peer device thread）強制 fence 前的存取排序在 fence 後之前，但不保證對其他 thread 可見——可見性需另靠volatile（繞過 L1）或 atomic。跨 block 規約常見模式：寫result[blockIdx.x]→ device fence →atomicInc(&count, gridDim.x)，最後完成的 block 拿到的舊值為gridDim.x - 1；少了 fence，counter 可能在部分和寫入前就遞增。__syncwarp(mask)預設 mask0xFFFFFFFF，只同步 mask 指名的 warp 成員。 - 回讀：同步原語與原子函式
Trap: Legacy atomic 只有 relaxed、不引入 fence；型別/CC 門檻與向量原子性陷阱
- 陷阱是什麼：以為
atomicAdd之類的 legacy 原子函式自帶記憶體排序；或以為對__half2/float4做 atomic 是整個向量單次原子；或對double/long long隨意挑運算。 - 為何容易混淆：「atomic」常被當成「順帶提供可見性與 fence」，型別支援也容易記成「都差不多」。
- 正確觀念：CUDA 有五種 atomic 提供方式，推薦 Extended
cuda::atomic/cuda::atomic_ref（效率/安全/可攜）。Legacyatomic<Op>()僅 device、記憶體序只有memory_order_relaxed、不引入 fence，只保證原子性；scope 由後綴定：無＝device、_block＝block、_system＝system。型別/CC 門檻要記：atomicAdd的__nv_bfloat16/162需 CC 8.x+、float2/float4需 CC 9.x+ 且僅 global 位址，double只有atomicAdd支援；atomicExch不支援double；unsigned short只有atomicCAS支援；帶號long long只有atomicMin/atomicMax；128-bit 模板（atomicExch/atomicCAS）需 CC 9.x+ 且 16-byte 對齊。向量型別逐元素原子，整個向量非單次原子。atomicInc(addr,val)＝old>=val?0:old+1、atomicCAS＝old==compare?val:old（回 old）。浮點 CAS 迴圈必須以整數位元樣式比較，因NaN!=NaN否則永不收斂而 hang。另有 Built-in__nv_atomic_*（CUDA 12.8+，scope 列舉 THREAD/BLOCK/CLUSTER/DEVICE/SYSTEM 共五級，order/scope 須整數字面量）。 - 回讀：同步原語與原子函式
Trap: mbarrier/pipeline primitives 的限制，與 grid.sync() 必須用 cooperative launch
- 陷阱是什麼：對
__pipeline_memcpy_async傳任意size_and_align；或以為this_grid().sync()像__syncthreads()一樣用<<<>>>啟動就能跨 grid 同步。 - 為何容易混淆：這些 device-callable 底層 API 介面相近，且
grid.sync()寫法和 block 同步幾乎一樣。 - 正確觀念：
__mbarrier_*（<cuda_awbarrier_primitives.h>）的bar必須位於__shared__；__mbarrier_arrive_and_drop比__mbarrier_arrive額外遞減下一 phase 的 expected count；test_wait立即查詢、try_wait可能掛起到 phase 完成或max_sleep_nanosec；重用該 shared 前須__mbarrier_inval。__pipeline_memcpy_async只追蹤 global→shared，size_and_align只能是 4/8/16（等於兩端指標對齊），zfillbytes 補零；__pipeline_wait_prior(N)等到批次L-N（含）。Cooperative Groups 三大 implicit group 由this_thread_block()/this_cluster()（CC 9.0+）/this_grid()取得；tiled_partition<Size>的Size須 2 的次方且 ≤ 1024，且shfl_up/shfl_down/shfl_xor/ballot僅 tile ≤ 32；cg::reduce硬體加速需 CC 8.0+ 且只有 4-byte 型別。grid.sync()必須用cudaLaunchCooperativeKernel（不可<<<>>>），僅 CC 6.0+，可先用cudaDeviceGetAttribute(..., cudaDevAttrCooperativeLaunch, dev)查支援。所有 collective 全員必須參與，否則 UB。 - 回讀：記憶體屏障、Pipeline 與 Cooperative Groups
Warp 函式與 CUDA 巨集
Trap: Volta+ warp 函式必須帶 mask（
_sync 版），且 __activemask() 不能拿來當 mask
- 陷阱是什麼：沿用舊式無 mask 的
__shfl/__any/__ballot；或用__activemask()的回傳值當下一個*_sync的 mask；或以為 warp 函式提供記憶體排序。 - 為何容易混淆：舊教材常見無 mask 版本，且
__activemask()「看起來」剛好回傳一組 active lane 可當參與集合。 - 正確觀念：vote/match/reduce/shuffle 與
__syncwarp都以*_sync結尾並帶 mask，mask 用來在硬體執行前確保收斂，效率最佳是全員參與（0xFFFFFFFF）——此為本節原文確認的事實。（以下為補充背景、非本節原文，源自 Independent Thread Scheduling／CUDA 9 棄用脈絡：Volta 起 thread 可獨立排程、不再保證 lock-step 收斂，故須以 mask 明確列出「應參與的 lane 集合」並強制收斂，這也是舊式無 mask__shfl/__any/__ballot被*_sync版取代的原因。）mask 鐵則：每個呼叫 thread 的 bit=1、非呼叫 thread 的 bit=0（退出 thread 被忽略），mask 內所有非退出 thread 須以相同 mask 執行（不同 mask 並發只在彼此 disjoint 時合法）。__activemask()只是瞬間快照、不代表「哪些 lane 走某分支」，且收斂在後續指令不保證延續——即使下一行就是*_sync，把__activemask()結果拿去當 mask 仍不安全（active 集合可能在兩步間改變），正解是直接用已知的靜態 mask。所有 warp*_syncintrinsic 不提供任何 memory ordering / barrier；shuffle 的width須為 2 的次方且 ∈[1,32]，讀到 inactive thread 的值未定義。__reduce_*_sync需 CC 8.x+。__CUDA_ARCH__只在 device code 有定義、值＝compute_<version>*10（compute_80→800）。 - 回讀：Warp 函式與巨集
Trap:
__trap() 損毀 CUDA context、DPX 三路 min/max+ReLU、__nanosleep 上限約 1 ms、位址空間/快取 intrinsic
- 陷阱是什麼：以為
__trap()只中止當前 kernel；不知 DPX 能一次做三路 min/max 並夾 ReLU；以為__nanosleep可睡任意久；或在跨 translation unit 的指標上仰賴編譯器自動判定位址空間。 - 為何容易混淆：這些 device 工具函式較少出現在課本，名字也不直覺，又常與 warp 同步原語混在一起記。
- 正確觀念：
__trap()從任一 device thread 呼叫即中止 kernel 並向 host 觸發中斷，導致 CUDA context 損毀、後續所有 CUDA 呼叫與 kernel 啟動皆失敗（可攜替代cuda::std::terminate()）；__brkpt()則只暫停 kernel 供除錯。__nanosleep(ns)約睡 ns 奈秒、上限約 1 ms，典型用於 exponential back-off mutex。DPX：對最多三個 16/32-bit 有/無號整數做 min/max 與 fused add-min/max，可選 ReLU（夾到 0）——如__vimax3_s32(a,b,c)＝max(a,b,c)、__vimax3_s32_relu＝max(a,b,c,0)、__viaddmax_s32＝max(a+b,c)；s16x2/u16x2後綴一次處理兩個 packed 16-bit，硬體加速與否視 CC，常用於 Smith-Waterman／Needleman-Wunsch／Floyd-Warshall 等動態規劃。位址空間判斷__isGlobal/__isShared/__isConstant/__isLocal（generic 指標指向該空間回 1，NULL 行為未指定）、轉換__cvta_generic_to_global（回size_t）/__cvta_global_to_generic（回void*）等對應 PTXcvta.*；低階快取載入/儲存__ldg（read-only L1/Tex 快取載入）、__ldcg/__ldcs、__stcg/__stwt依 PTX 快取運算子操作（x3 向量分量除外）。 - 回讀：Warp 函式與巨集
編譯器提示與 Warp Matrix
Trap: assert 觸發 context 損毀、
__builtin_assume(false) 是 UB、診斷只有 warning 可被抑制
- 陷阱是什麼：以為 device
assert()失敗只中止當前 kernel、之後還能繼續用 CUDA；把__builtin_assume(pred)當成執行期檢查；以為nv_diag_*能把任何診斷都抑制掉。 - 為何容易混淆：assert 在 host 端只是中止單一程序；
assume/expect/unreachable名字像普通檢查或標準屬性；diagnostic pragma 又長得像可隨意開關。 - 正確觀念：device
assert(expr)在expr==0時中止 kernel，經cudaDeviceSynchronize/cudaStreamSynchronize/cudaEventSynchronize同步後向 stderr 印 block/thread 與訊息，並使 CUDA context 損毀——之後任何 CUDA 呼叫或 kernel 啟動都失敗並回cudaErrorAssert；production 用NDEBUG（include 前定義或-DNDEBUG）停用，且expr不應有副作用（停用會改變功能），assert在__device__/__tile__皆可用。最佳化提示：__builtin_assume(pred)讓編譯器假設 pred 為真，執行期為 false 是 UB、有副作用是 unspecified；__builtin_expect(in,exp)回傳in並提供分支預測（類[[likely]]/[[unlikely]]）；__builtin_unreachable()告知控制流不會到達、到達即 UB；__builtin_assume_aligned(ptr,align)的align須為 2 的次方整數字面值；__assume()僅 MSVC。診斷 pragma（nv_diag_suppress/nv_diag_warning/nv_diag_error/nv_diag_default，另有nv_diagnostic push/pop）：任何診斷皆可升成 error，但只有 warning 能被抑制、或在升 error 後再還原嚴重度；只影響 nvcc CUDA 前端、不影響 host 編譯器（支援時定義__NVCC_DIAG_PRAGMA_SUPPORT__）。#pragma nv_abi（separate compilation 保留暫存器；放 call site 只對 free function 的間接呼叫有效）與#pragma nv_mma_throughput（僅 entry-function、實驗性）本身不引發 trap。 - 回讀：編譯器提示與 Warp Matrix
Trap: WMMA fragment 是架構特定的不透明結構，不可跨介面互傳
- 陷阱是什麼：把
wmma::fragment當普通結構在函式間／跨編譯單元傳遞；或假設其元素映射、暫存器配置固定；或忽略對齊與全 warp 一致呼叫的要求。 - 為何容易混淆：fragment 看起來就是個一般 C++ 模板物件，能宣告就以為能自由傳。
- 正確觀念：WMMA（
nvcuda::wmma）用 Tensor Core 做D=A*B+C，需 compute capability 7.0+（bfloat16/tf32/double/bmmaBitOpAND需 8.0、bmmaBitOpXOR需 compute_75+）、整個 warp 協作，且只有在條件對整個 warp 求值相同時才能用於條件式碼，否則極可能 hang。fragment 是不透明、架構特定的 ABI 結構，元素到內部儲存的映射未指定且隨架構改變；若兩函式為不同 link-compatible 架構（如 sm_70 與 sm_75）編譯卻連結在一起並互傳 fragment，會導致錯誤結果或損毀，且編譯期/執行期工具可能都偵測不到——應先store_matrix_sync把矩陣存到記憶體，再以指標安全傳遞。對齊與一致性：mptr須 256-bit 對齊，ldm為__half8 的倍數 /float4 的倍數（皆 16 bytes 倍數）；mptr/ldm/layout與所有模板參數對全 warp 須一致且由全員呼叫。fragment的Use：matrix_a(m×k)、matrix_b(k×n)、accumulator(m×n)；double尺寸 8x8x4 以.rn捨入、tf32 唯一尺寸 16x16x8、satf=true時 NaN 結果存+0。另記#pragma unroll：無引數＝完全展開、0/1＝不展開、負值/超INT_MAX＝忽略並警告。 - 回讀：編譯器提示與 Warp Matrix
浮點數與數學函式
Trap: GPU 結果和 CPU 不同不代表算錯——FMA contraction、結合律、denormal/-ftz 的影響
- 陷阱是什麼：看到 GPU 與 CPU 浮點結果位元不同就斷定 GPU 有 bug；或以為 IEEE-754 相容就代表能動態切 rounding mode、會偵測例外。
- 為何容易混淆：「都符合 IEEE-754」直覺以為結果應一模一樣。
- 正確觀念：差異常源自 rounding、FMA contraction、加乘不滿足結合律、編譯器優化、library 實作差異，並非算錯。CUDA 遵循 IEEE 754-2019，但無動態可設 rounding mode（四種 rn/rz/ru/rd 須以特定命名的 device intrinsic 對個別運算選，預設 rn）、無浮點例外偵測（如同永遠 masked，SNaN 不 signaling）。FMA
a*b+c只 round 一次並保留 double-width 乘積防 subtractive cancellation，比分兩步更精確；x*y+z在-fmad=true（含於--use_fast_math）時 contract 成 FMA。denormal（subnormal）較慢，-ftz=true（flush-to-zero，含於--use_fast_math）可沖成 0。最佳可移植性用 nvcc 預設-ftz=false、-prec-div=true、-prec-sqrt=true且不開--use_fast_math；但 re-association 與 FMA contraction 預設仍允許（等同--fmad=true），要關須另設-fmad=false。+ - * / 1/x在 rn 下 0 ULP；-prec-div=false後/→2 ULP、1/x→1 ULP。global memory 上 atomic single-precision add 永遠 flush-to-zero、shared memory 上永遠支援 denormal（不論-ftz）。浮點除以 0 是 implementation-specific（得 infinity）不是 UB，整數除 0 才是 UB；NaN==NaN為 false。 - 回讀：浮點數運算基礎
Trap: standard math 函式 vs intrinsic（
__sinf 等）的精度/速度取捨與 --use_fast_math
- 陷阱是什麼：分不清
sinf與__sinf、fdividef與__fdividef；或全域開--use_fast_math卻沒意識到精度與特例處理被改變。 - 為何容易混淆：名字只差兩個底線，又都「算同一個數學函式」。
- 正確觀念：兩條路徑——標準路徑（
cuda::std::，header<cuda/std/cmath>，host+device，有 ULP 界）與 intrinsic 路徑（__前綴、只在 device code、映射較少原生指令、較快但較不精確）。--use_fast_math會把一組 device Math API（含cuda::std）自動換成 intrinsic：x/y/fdividef→__fdividef、sinf/cosf/tanf→__sinf/__cosf/__tanf、logf/expf/powf/tanhf等同理，且sqrtf從 0 ULP 變 1 ULP。fdividef≠__fdividef：非標準的fdividef(x,y)是 0 ULP 等同x/y，而 intrinsic__fdividef(x,y)對|y|∈[2⁻¹²⁶,2¹²⁶]是 2 ULP。intrinsic 函式不受-prec-div/-prec-sqrt/-fmad影響，唯-ftz=true例外。捨入後綴_rn/_rz/_ru/_rd＝最近偶數/趨零/向上/向下；單雙精度基本 intrinsic（__fadd_/__fmul_/__fmaf_…）為 0 ULP 且 IEEE 相容，且__fadd_/__dadd_/__fmul_/__dmul_絕不會被合併成 FFMA/DFMA（而+、*常被合併）。取整用rint/rintf（單指令、ties to even），勿用round/roundf（多指令、ties away from zero）。Max ULP 由測試得出、不保證。另有一批 Non-Standard 擴充函式（rsqrtf/rsqrt、norm3d(f)/norm4d(f)、Besselj0f/y0f/cyl_bessel_i0f等；對__half/__nv_bfloat16/__float128只在 device code），其中norm(dim,p)/rnorm(dim,p)採快速演算法而無誤差界。整數位元 intrinsic（__popc計 set bit、__ffsfind-first-set、__clzcount-leading-zeros、__brevbit-reverse）屬整數而非浮點，誤差/規格另查對應章節。 - 回讀：數學函式與 Intrinsic
裝置端 Runtime
Trap: device runtime 只是 host runtime 的子集——CDP2 沒有 cudaDeviceSynchronize/cudaStreamSynchronize
- 陷阱是什麼：在 kernel 內用
cudaStreamCreate()、cudaStreamSynchronize()/cudaStreamQuery()、（CDP2）cudaDeviceSynchronize()；或把 device 端建立的 stream/local/shared 指標傳給 child；或以為 launch 後沒錯誤就代表 child 成功完成。 - 為何容易混淆：device runtime 與 host CUDA Runtime API 語法相同，自然以為功能也一樣。
- 正確觀念：CUDA device runtime（kernel 內可呼叫，最常用於 Dynamic Parallelism (CDP2) 與 Device Graph Launch，原型編譯時自動引入）只是 host runtime 的子集：無多 GPU（只能操作當前 device，但可查任意 device 屬性）、stream 必須用
cudaStreamCreateWithFlags(cudaStreamNonBlocking)（cudaStreamCreate()不可用）、event 須cudaEventDisableTiming、memcpy 只 async / 只 device-to-device / 不可傳 local/shared 指標。未列於支援子集的有cudaStreamSynchronize/cudaStreamQuery/cudaDeviceSynchronize（CDP2）。要確認子 kernel 完成改用cudaStreamTailLaunch（多數情況等同cudaDeviceSynchronize效果），另有cudaStreamFireAndForget（立即排程、不被相依）；兩者都需 64-bit 編譯且不可定義CUDA_FORCE_CDP1_IF_SUPPORTED，且不能 record/wait event（否則cudaErrorInvalidValue）。stream handle 為建立它的 grid 私有，不可傳給 parent/child。pending launch slot 預設 2048、event slot 為 limit 兩倍；cudaGetLastError是 per-thread；launch 無錯誤不代表 child 成功完成；__constant__device 端唯讀，且因 symbol 可直接&取址，cudaMemcpyToSymbol/cudaGetSymbolAddress不需要也不支援。%smid/%warpid是 volatile、不可假設不變。 - 回讀：CUDA 裝置端 Runtime
記憶體模型與執行模型
Trap: atomic 只在「指定的 scope」內 atomic——scope 不涵蓋對方就構成 data race (UB)
- 陷阱是什麼：以為只要用了 atomic 操作就一定 race-free；message passing 中 store 端用
thread_scope_block、load 端在別的 block。 - 為何容易混淆：「atomic」直覺就是「安全」，忽略了 thread scope 是否涵蓋對方執行緒。
- 正確觀念：標準 C++ 假設同步成本均勻，CUDA C++ 不同——距離越遠成本越高，故在
cuda::namespace 以 thread scope 擴充。四級由大到小：thread_scope_system（含多 GPU/CPU）→thread_scope_device（同一 device 且同一 memory synchronization domain）→thread_scope_block（同一 block）→thread_scope_thread（只與自己）。一個 atomic 操作只在它指定的 scope 下保證為 atomic；scope 不是 system 時自動成立，是 system 時須滿足裝置屬性（system allocated memory 需pageableMemoryAccess、managed 需concurrentManagedAccess、mapped 需hostNativeAtomicSupported）。data race（含 scope 版）：兩個可能並行的衝突動作，至少一個在「涵蓋對方執行緒的 scope」下不是 atomic，且彼此無 happens-before，即構成 data race → UB。經典 message passing：store 用 block scope（由 Block 0 的 thread 執行、只涵蓋 Block 0），load 的 thread 在 Block 1 不被涵蓋，於是兩者對彼此不是 atomic＝race；正解是兩端都用thread_scope_device。std::/cuda::std::型別等同cuda::型別以thread_scope_system實例化。 - 回讀：CUDA C++ 記憶體模型
Trap: device thread 的 forward progress 要靠 API/依賴觸發——spin-loop 內要反覆 query
- 陷阱是什麼：以為 launch 後 device thread 必然會自己開始跑；或在 spin-loop 之外只呼叫一次
cudaStreamQuery/cudaDeviceSynchronize就期待死結解開；或兩個互相等待的 kernel 放在不同 stream。 - 為何容易混淆：host thread 的 concurrent FP 直覺被誤套到 device thread。
- 正確觀念：CUDA 對 device thread 提供 parallel forward progress [intro.progress.9]（前提是 host 提供 concurrent FP）：thread 尚未執行任何步驟前不保證會開始，一旦執行一步才升級為 concurrent FP。傳染範圍：某 thread 推進後，若屬 Cooperative Grid → 整個 grid、否則 → 整個 thread-block cluster（蘊含整個 block）；其他 cluster 不保證推進。device thread 比 host 少了
std::this_thread::yield與 trivial infinite loop 兩種「最終會做」的選項，且對 automatic storage（區域變數）做 volatile/atomic 不算進度、只認 atomic read（單純 atomic write/fence 不算，刻意為效能）——故while(true){}、while(true) yield();、對區域volatile/atomic spin 都可能無人推進。任一 CUDA API 呼叫最終必 return 或確保至少一個 device thread 推進（推進者不必與該呼叫相關），且 query 函式不得在無 thread 推進時持續回cudaErrorNotReady——所以單次 query 不保證 device 啟動，唯有 spin-loop 內反覆 query（或一定推進整批工作的 sync API）才保證最終推進。device thread 在所有依賴完成前不得啟動：跨不同 stream 無依賴可能 starvation，同一 stream 的順序依賴才確保依序執行。粗測工具：CUDA_DEVICE_MAX_CONNECTIONS=1 CUDA_LAUNCH_BLOCKING=1（非充分）。 - 回讀：CUDA C++ 執行模型