面試陷阱題 (Exam Traps)
Purpose
本頁彙整 PMPP 全書最高頻、最易考、最容易答錯的「陷阱題」。每則拆成 誤解 → 為何錯 → 正解,並連回原始概念筆記。複習前掃過這份清單,能快速找出自己的弱點與盲點;面試 / 考前各章只看一兩則最關鍵的即可。
基礎與架構 (Foundations & Architecture)
陷阱:把 parallel 部分加速到無限大就能無限加速整體
- 誤解:只要把可平行部分加速到極致,整體 speedup 就能無限放大
- 為何錯:Amdahl's Law 中 serial 部分 (1-p) 不會改變,構成整體加速的硬上限 1/(1-p);p=30% 時整體上限僅 1.43×
- 正解:整體 speedup 由 sequential fraction 封頂;必須讓 >99.9% 執行時間落在 parallel 部分才可能 >100×
- Amdahl's Law
陷阱:認為 GPU 在所有任務上都比 CPU 快
- 誤解:GPU 核心多、算力高,任何工作都該丟給 GPU
- 為何錯:GPU 是 throughput-oriented;降低延遲遠比提高吞吐昂貴,對只有少數 thread 的序列工作,低延遲的 CPU 反而快很多
- 正解:序列 / 少 thread 的部分跑 CPU,大量 thread 的數值密集部分跑 GPU(CPU+GPU 異質協同)
- 異質協同運算
陷阱:把 cudaMemcpy 當英文「copy A to B」寫成 (src, dst, ...)
- 誤解:依照「把 A 複製到 B」的語感,把來源指標放第一個參數
- 為何錯:cudaMemcpy 採 C memcpy 順序 (dst, src, size, kind);寫反會靜默複製錯方向,不報錯卻得到垃圾值
- 正解:目的地指標在前 — cudaMemcpy(A_d, A_h, size, cudaMemcpyHostToDevice);cudaMalloc 第一參數則須 (void*)&A_d,size 用 nsizeof(float)
- Device Memory 與資料傳輸
陷阱:省略 if(i < n) 邊界檢查,以為 grid thread 數剛好等於 n
- 誤解:啟動的 thread 數正好等於資料量 n,不必檢查邊界
- 為何錯:grid thread 數 = ceil(n/blockDim.x)*blockDim.x,幾乎總是大於 n;尾端多餘 thread 會越界讀寫 global memory
- 正解:kernel 內務必加 if(i < n),只有 i < n 的 thread 做事,如此同一 kernel 才能處理任意長度
- Kernel Functions 與 Threading
陷阱:認為 grid 三個維度上限都是 2^31-1
- 誤解:gridDim 的 x/y/z 三維上限相同
- 為何錯:只有 gridDim.x 上限是 2^31-1;gridDim.y 與 gridDim.z 上限只到 65535 (2^16-1)
- 正解:x:1..2^31-1;y,z:1..65535。另注意 block 總 thread 數須 xyz <= 1024,(32,32,2)=2048 非法
- 多維 Grid 組織
陷阱:把 thread 的 .x 對到 row、.y 對到 col
- 誤解:直覺把第一個維度 .x 當成資料的第一個維度 row
- 為何錯:grid/block 維度排序是 x 先,而資料 / 數學排序是 row(=y)先,兩者剛好相反
- 正解:.y ↔ row(垂直)、.x ↔ col(水平);row 用 blockIdx.y/threadIdx.y,col 用 .x
- 矩陣乘法 Kernel
陷阱:在 if-then-else 每條 branch 各放一個 __syncthreads() 就算合法同步
- 誤解:只要每條路徑都有 barrier,block 內 thread 就都同步了
- 為何錯:then-path 與 else-path 的兩個 __syncthreads() 是兩個不同的 barrier;thread 被分到不同 barrier,沒有任何一個被全體到達 → deadlock / undefined behavior
- 正解:__syncthreads() 必須被 block 內每個 thread 執行;同步點不可放在會讓部分 thread 跳過的分支裡
- Block 排程與同步
陷阱:占用率 (occupancy) 越高就一定越快,多用一點 register 只會小幅下降
- 誤解:把 occupancy 當成優化的最終目標,且資源用量對它是連續影響
- 為何錯:occupancy 只是提升 latency tolerance 的手段;資源以 block 為單位整批分配,跨過邊界 (31→33 regs) 會使 blocks/SM 驟降 → performance cliff
- 正解:occupancy 階梯式掉落,須與記憶體頻寬、coalescing 等真正瓶頸一起評估,不是優化終點
- Resource Partitioning 與 Occupancy
記憶體與效能 (Memory & Performance)
陷阱:把 local memory 當成 on-chip 高速記憶體
- 誤解:名稱叫 "local",以為靠近 thread、速度快
- 為何錯:local memory 實際位於 off-chip global memory,延遲與 global 相同,只是 thread 私有
- 正解:local memory 存放自動陣列、spilled registers、call stack,速度等同 global;只有自動「純量」才進 register
- CUDA 記憶體類型
陷阱:認為 tiling 把 global memory 存取減少 TILE_WIDTH² 倍
- 誤解:block 是 TILE_WIDTH×TILE_WIDTH 二維,所以減量也是平方
- 為何錯:減量只與 tile 的單一維度成正比,不是面積
- 正解:減量因子 = TILE_WIDTH(16×16 tile 減 16 倍);此外載入後與使用後的兩個 __syncthreads() 分別防 RAW 與 WAR,缺一即 race
- Tiling 與 Tiled 矩陣乘法
陷阱:把 coalescing 當成「同一 warp 存取同一位址就會合併」
- 誤解:以為 warp 內 thread 讀同一個位址才會被合併
- 為何錯:coalescing 的條件是同時存取「相鄰連續位址」(thread i → X+i),不是同一位址,且需 warp 的 SIMD 同步排程
- 正解:連續 threadIdx.x 對應連續位址 (stride 1) 才合併成一個 burst;column-major 輸入可用 corner turning 經 shared memory 中轉
- Memory Coalescing 與 DRAM Bursts
陷阱:認為 thread coarsening 對所有 kernel 都能加速,且 COARSE_FACTOR 越大越好
- 誤解:把粗化當成萬用優化,因子越大省越多
- 為何錯:只有存在「平行的代價」(重複載入 / 計算 / 同步) 且硬體會串行化 block 時才有效;vector add、RGB→grayscale 沒有代價,套用無益;因子過大會餓死硬體並壓低 occupancy
- 正解:先確認有可省的重複成本;最佳 COARSE_FACTOR 是 device-specific 與 dataset-specific,須調校
- Thread Coarsening
平行模式 (Parallel Patterns)
陷阱:認為 "convolution kernel" 指 CUDA kernel function
- 誤解:把卷積的 "kernel" 與 CUDA 的 global kernel 混為一談
- 為何錯:PMPP 為避免命名衝突,把卷積的權重陣列改稱 filter (F)
- 正解:convolution kernel = filter(權重陣列),與 CUDA kernel function 無關;基本卷積真正瓶頸是 memory bandwidth(arithmetic intensity 僅 0.25 OP/B),不是 control divergence
- Convolution 基礎與基本 Kernel
陷阱:把 halo cells 和 ghost cells 混為一談
- 誤解:兩者都是 tile 周圍的「額外」cell,所以歸為同類
- 為何錯:角色完全不同,且 tiled kernel 的兩個 if 也各管不同事
- 正解:halo 是鄰近 tile 的真實內部元素(有效資料);ghost 是陣列邊界外不存在的元素(填 0);input tile 變大是因 halo 而非 ghost
- Tiled Convolution 與 Halo 處理
陷阱:認為 stencil order 等於總格點數或維度
- 誤解:3-point stencil 被誤判為 order 3,3D 7-point 被誤判為 order 7
- 為何錯:order 指的是中心點「每一側」的格點數 = 近似導數的階數
- 正解:3-point = order 1、3D 7-point = order 1(每軸每側 1 點);stencil pattern 稀疏,tiling 重用率本質遠低於 dense convolution
- Stencil 基礎與基本核心
陷阱:認為 atomicAdd 會強制 thread 之間特定的執行順序
- 誤解:atomic 會讓 thread 依某種順序依序更新
- 為何錯:atomic 只保證同位址 RMW 互斥(序列化),哪個 thread 先到是任意的;結果正確只因加法可交換
- 正解:atomicAdd 序列化但不排序;且同位址高度競爭時 throughput 崩到 ~1/(2×latency),由 latency 而非 bandwidth 主導;呼叫時位址在前 — atomicAdd(&Total, Partial)
- 原子操作與基本直方圖
陷阱:認為 reduction tree 比 sequential 做更少的總運算
- 誤解:平行樹「同時做」,所以總運算量比循序少
- 為何錯:tree 與 sequential 的總運算數都是 N-1;平行化只縮短步數,不減少 work
- 正解:reduction tree 是 work-efficient(仍做 N-1 次運算),換到的是時間步數從 N 降到 log2(N)
- Reduction 基礎與簡單 Kernel
陷阱:Brent-Kung 因為 work-efficient 所以一定比 Kogge-Stone 快
- 誤解:總運算量少 = 比較快
- 為何錯:work efficiency 高不代表步數少;在 execution units 無限時,Brent-Kung 因多了反向樹階段需約 2× 步數,反而更慢
- 正解:Brent-Kung 只在執行單元 P 有限時贏(總運算量主導);資源充足時 Kogge-Stone 贏(步數主導)
- Brent-Kung 與 Thread Coarsening
也可以
- 誤解:相等時用嚴格小於 < 比較沒差
- 為何錯:嚴格小於會在 key 相等時取 B,變成 B 優先,破壞 across-list stability(穩定排序)
- 正解:tie 必須用 A[i] <= B[j] 讓 A 元素優先;且每個 thread 的 input 範圍須由 co-rank(binary search)反推,非固定 index 算式
- Merge 基礎與 Co-rank
陷阱:認為 radix sort 也受 O(N log N) 下界限制
- 誤解:所有排序都不可能優於 O(N log N)
- 為何錯:O(N log N) 只是 comparison-based 的理論下界;radix sort 是 noncomparison-based,靠 digit 分桶可更快
- 正解:radix 可優於 O(N log N),代價是只適用 integer 等特定 key;且 iteration 間有序列依賴,只在單一 iteration 內 (one-thread-per-key) 平行
- 排序基礎與平行 Radix Sort
應用 (Applications)
陷阱:認為 CSR 一定比 COO 好,所以總是選 CSR
- 誤解:CSR 較精簡,任何情況都該用 CSR
- 為何錯:CSR 跨 row 平行在 row 數太少時餵不飽 GPU;且連續 thread 各負責不同 row,讀各 row 起點彼此相距很遠 → 非 coalesced,又有 control divergence
- 正解:row 數少或需動態建構 / 修改時 COO 更佳;COO 才 coalesced(連續 thread 讀連續元素),但同 row 共享輸出需 atomicAdd,CSR 則不需
- COO 與 CSR 格式
陷阱:認為光靠 CSR 就能做最短路徑回溯
- 誤解:有了 CSR 的鄰接資訊就能找前驅、回溯路徑
- 為何錯:CSR(srcPtrs+dst)只給 outgoing edges,無法找 predecessor
- 正解:路徑回溯要用 CSC(dstPtrs+src)取得每個 vertex 的 predecessors,挑 level 比自己小 1 的;另 BFS 的 level 完全取決於 root
- 圖的表示法與 BFS
陷阱:認為 BFS push kernel 標記 neighbor 也需要 atomic
- 誤解:多 thread 可能同時寫同一 neighbor,所以要 atomic 保護
- 為何錯:多個 thread 用「相同的 currLevel」標記同一未訪問 neighbor 是 idempotent 寫入,不是 read-modify-write,結果與次數無關
- 正解:無 frontier 的 push 不需 atomic;只有改用 frontier(避免重複入隊浪費)時,才需 atomicCAS 確保只有一個 thread 入隊
- Vertex / Edge-centric BFS
陷阱:認為 conv layer 的 7 層迴圈都能無腦平行化
- 誤解:7 層迴圈彼此獨立,全部開成 thread 即可
- 為何錯:內層 c/p/q 會累加進同一個 Y output element,造成 read-modify-write race,需 atomic 才能平行
- 正解:只有 n/m/h/w 4 層是 easy parallel;c/p/q 預設保持序列。另外 im2col/GEMM 並不一定省記憶體(輸入複製最多 K·K 倍),靠 cuDNN 的 on-chip lazy materialization 才真正解決膨脹
- GPU Conv Layer 與 GEMM
陷阱:以為 FHD 的 scatter atomics 可用 shared-memory privatization 解決(像直方圖)
- 誤解:照搬 histogram 的 privatization,把輸出私有化到 shared memory
- 為何錯:輸出陣列 rFhD/iFhD 大小 = 總 voxel 數(數百萬),放不進 shared memory,privatization 不可行
- 正解:改 scatter 為 gather(先 loop fission 再 interchange),每 thread 擁有唯一輸出 voxel,徹底消除寫入衝突與 atomic
- FHD Kernel 平行化結構
陷阱:認為應該平行化「最佳化過」的序列程式碼
- 誤解:既然要加速,就拿最佳化後的版本來平行化
- 為何錯:最佳化後 atom 在最外層 → 變成 scatter,多 thread 寫同一 grid point 需 atomicAdd,嚴重序列化
- 正解:首選 gather — 平行化「未最佳化」版本,每 thread 擁有一個 grid point、無需 atomic;atoms[] 放 constant memory 由 constant cache broadcast,不受 global memory bandwidth 限制
- DCS Scatter vs Gather
進階實務 (Advanced Practices)
陷阱:整體 speedup 由加速最猛的平行段決定,序列段只佔 5% 可忽略
- 誤解:把心力全放在加速平行段,小小的序列段不重要
- 為何錯:平行段加速 100× 也壓不動序列段;5% 序列段使上限只到 20×(Ch.17 的 CG solver 僅佔 1% 卻成 speedup 限制因素)
- 正解:整體 speedup 被 sequential fraction (1-p) 封頂;小序列段累積後是真正瓶頸,須用 task-level / hierarchical MPI-CUDA 壓低它
- Amdahl's Law 與目標
陷阱:把 SIMT、SPMD、SIMD 當成同義詞
- 誤解:三個縮寫指同一件事
- 為何錯:三者處於不同層級,混用會答錯執行模型題
- 正解:SIMD 是硬體 vector lane;SIMT 是 GPU warp 鎖步但各 thread 有自己狀態;SPMD 是同一份程式跑不同資料的程式設計模型(含 MPI ranks)
- 計算思維與平行化策略
陷阱:MPI collective(MPI_Bcast / MPI_Barrier)只需用到的 process 呼叫
- 誤解:只有需要該資料的 process 才呼叫 collective
- 為何錯:collective 定義在整個 communicator 上,任一 process 沒到達同一個 collective 呼叫就 deadlock
- 正解:communicator 內所有 process 都必須呼叫該 collective;勿用 per-rank 條件排除部分 rank。手寫 send/recv 通常也輸給高度優化的 collective
- Collective 與 CUDA-aware MPI
陷阱:cudaMemcpyAsync 可搭配任何 malloc 的 host 指標
- 誤解:Async 版只要傳 host buffer 就會非同步
- 為何錯:async copy 需 page-locked(pinned)memory;pageable buffer 會退化成額外 staged copy 且變同步,async 語意悄悄失效
- 正解:用 cudaHostAlloc(..., cudaHostAllocDefault) 配置 pinned host buffer,以 cudaFreeHost 釋放;注意 boundary slice(送出)與 halo slice(收入)角色不同
- 計算與通訊重疊
陷阱:N_LINES=1024、BLOCK_DIM=64 時 launch 16 個 child kernel
- 誤解:把 child kernel 數當成 parent block 數 (1024/64=16)
- 為何錯:16 只是 parent block 數;child grid 由每個 parent thread 發起,不是每 block 一個
- 正解:共 launch 1024 個 child kernel(一 thread 一個);且省略 stream 時同 block 內所有 launch 共用 NULL stream 被序列化,需 per-thread named stream 才能並行
- Bezier Curves 範例
陷阱:UVA 讓 host code 可直接 dereference 任何 cudaMalloc 的 device 指標
- 誤解:有了統一虛擬位址空間,host 就能直接讀寫 device 指標
- 為何錯:UVA 只保證單一虛擬位址空間(故 cudaMemcpy 方向可自動偵測),不保證可存取性
- 正解:host 仍不能 deref cudaMalloc 指標;跨側存取任意結構需 Unified Memory + page fault(Pascal);zero-copy 走 PCIe(<10% 頻寬)只適合偶發稀疏存取,且 managed memory 仍需自行用 barrier / atomic 處理併發
- Host/Device 互動模型
陷阱:認為一般 grid 也能呼叫 grid.sync() 做跨 block 同步
- 誤解:既然有 grid-wide barrier API,任何 grid 都能用它跨 block 同步
- 為何錯:一般 grid 採 transparent scalability,block 數可超過 SM 容量而分波執行;全 grid barrier 會讓尚未上場的 block 永不啟動 → deadlock
- 正解:跨 block 同步須用 cooperative kernel(block 數不超過 GPU 容量,runtime 保證全部併發),以 cudaLaunchCooperativeKernel() 啟動
- Kernel 執行控制
陷阱:認為 GPU 進步是 compute 與 bandwidth 同步成長,memory 優化已不重要
- 誤解:新世代 GPU 又快又有大頻寬,tiling / coalescing 可以少做
- 為何錯:G80→A100 compute ×452 但 bandwidth 只 ×18,兩者差約 25×
- 正解:落差代表 arithmetic intensity 需求升高、memory wall 加深,tiling / coalescing 等優化反而更關鍵;最高潛力區是 storage-access parallelism
- 未來展望
數值 (Numerical)
陷阱:把 exponent bias 算成 2^e 或 2^(e-1)
- 誤解:偏移量直接用 2^e 或少減 1
- 為何錯:算錯會讓整個指數偏移、所有還原數值都錯;且若改用 two's complement 存指數,符號位會破壞無號比較的單調性
- 正解:bias = 2^(e-1) - 1(8-bit E → 127、11-bit → 1023);用 excess/biased 編碼,unsigned comparator 即可正確比較有號指數
- 浮點數資料表示
陷阱:把 precision 與 accuracy 當成同一回事
- 誤解:精度和準確度可以互換使用
- 為何錯:precision 是「靜態」的表示能力(看 mantissa 位數),accuracy 是「動態」的運算引入誤差(看運算如何實作),考試最常混淆這兩者
- 正解:precision 由 mantissa 位數決定(每多 1 bit 最大表示誤差減半);accuracy 由執行的運算(rounding / 近似演算法)決定,與 mantissa 位數無關
- 可表示數值、精度與準確度
陷阱:以為所有運算(含除法、sin)的捨入誤差都 ≤ 0.5 ULP
- 誤解:正確捨入保證任何運算誤差都在半個 ULP 內
- 為何錯:只有正確捨入的 add/sub ≤ 0.5 ULP;division 與超越函數 (sin, 1/x) 以 polynomial approximation 實作,項數不足時誤差可 > 0.5 ULP(早期 inversion 達 2 ULP),而 round-to-zero 截斷可達 1 ULP
- 正解:add/sub ≤ 0.5 ULP,division / 超越函數可能更大;使用 __sinf 等 hardware intrinsics 須權衡精度損失
- 可表示數值、精度與準確度
陷阱:假設「平行 reduction = 更準」或「循序 = 更準」
- 誤解:認定某一種加總順序在數值上一定比較準
- 為何錯:浮點加法不符結合律,(a+b)+c ≠ a+(b+c);平行樹與循序加總結果可能不同,課本範例平行較準只是巧合;大數加極小數時小數對齊右移會被整個吞掉 (swamping)
- 正解:不要假設方向;用 presort 升冪 + 相近 magnitude 分組 / Kahan 補償加總提升準確度,且排序方向要配合分組策略
- 演算法考量與數值穩定性
陷阱:以為可解的線性系統用 naive Gaussian elimination 一定能解出
- 誤解:只要解存在,直接做高斯消去就能求出
- 為何錯:若 lead 係數為 0,正規化步驟會除以零而失敗,即使解存在仍解不出 → numerically unstable
- 正解:用 pivoting — 找 lead 係數絕對值最大的 row swap 到最上方;多 block / cluster 的全域檢查昂貴時改用 partial pivoting / randomization(communication-avoiding)
- 演算法考量與數值穩定性