進階實務與未來演進 練習題 (Practice - Model of Host/Device Interaction and Memory Evolution)

(1) I/O 效率:disk controller、NIC 等 I/O 裝置只在 host memory 上運作；資料須先 H2D 搬入運算、運算完再 D2H 搬出才能做 I/O，增加 I/O latency、降低 throughput。(2) 大型資料結構:早期 GPU device memory 很小，逼開發者把大資料結構切塊 (partition) 才塞得下(如 Ch.18 把 3D 能量網格切成 2D slices)；有些應用甚至沒有好的切法。根源是早期 GPU 的記憶體架構限制，非 CUDA 語言本身。

用 cudaHostAlloc(..., cudaHostAllocMapped) 配置 pinned host memory(pinned 才不會被 OS page out)。host 指標須先經 cudaHostGetDevicePointer() 換成 device 指標再傳給 kernel；kernel deref 時走 PCIe / system interconnect 直接讀寫 host memory，免 cudaMemcpy()。但 system interconnect 頻寬 < global memory 的 10%，故只適合偶爾、稀疏存取的資料結構。

UVA 把 host 與 device 的獨立虛擬位址空間合併成單一共用虛擬位址空間 (UVAS)，保證每個實體位址只映射到唯一虛擬位址。runtime 只看指標值即可判斷它指向 host 還是 device → cudaMemcpy() 不再需要指定方向(可用 cudaMemcpyDefault)。但 UVA 不保證 accessibility:host code 仍不能直接 deref cudaMalloc() 回傳的 device 指標(反之亦然)；唯一例外是 zero-copy 指標可直接當 kernel 參數。

早期 GPU 用 32-bit 虛擬 + 最多 32-bit 實體位址 → 可定址 $2^{32}$ = 4 GB，device memory 與 kernel 操作的資料集都不能超過 4 GB；且 32-bit VA 裝不下現代 CPU 的 48-bit VA。Kepler 起改用 64-bit 虛擬位址 + ≥40-bit 實體位址 → 可裝 >4 GB DRAM、host/device 可用完全相同的指標值指同一資料，並把多顆 GPU 的 device memory 併入統一實體位址空間 → 一顆 GPU 可直接 deref 指標存取另一顆 GPU 的記憶體(peer access，免 host 觸發的 D2D copy)。

用 cudaMallocManaged() 配置 CPU/GPU 共享的 managed memory pool，單一指標存取，runtime + 硬體負責 migration/coherence。Kepler/Maxwell 上的三限制:(1) 每次 grid launch 前，CPU 改過的 managed memory 全部要 flush 到 GPU；(2) CPU 與 GPU 不能同時存取同一塊 managed allocation；(3) unified address space 受限於 GPU 實體記憶體大小。根本原因:這些世代缺乏 host/device 間的硬體 coherence，migration 多由軟體完成。

兩大特性:(1) 49-bit 虛擬位址(足以涵蓋現代 CPU 的 48-bit VA + GPU 自身記憶體) → 整個系統(所有 CPU/GPU)成為單一虛擬位址空間；(2) memory page fault handling。有了 page fault，CUDA runtime 改用 invalidate-based coherence:host/device 修改 managed 變數時用 page mapping/protection 讓對方副本失效；launch 時不必把所有 GPU 副本更新到最新——若 grid 存取到被 invalidate 或不在 device 的 page，就觸發 page fault，按需 migrate 到 GPU(或 map-on-access 走 interconnect)後恢復執行。

早期 compiler 必須把所有函式 body inline 進 kernel object，runtime 沒有真正的 call，故不支援 system call、dynamic library call、recursion、virtual function。Kepler/CUDA 5 起由 per-thread 的 massively parallel call frame stack(有 cache、快速實作)支撐真正的 runtime call，compiler 不再被強制 inline。效益:composability(不同作者寫不同元件再組裝、廠商可釋出不含原始碼的 device library 做 IP 保護)、可呼叫 printf()/malloc()(printf() 對 production 除錯特別有用，讓非技術 end user dump 內部狀態回報 bug)。

解決 irregular data 的 load imbalance、以及多 block 想用 mutex 等互斥機制合作(如共享 work queue)時的 deadlock 風險。它保證所有 thread block 同時併發執行(最多到填滿 GPU 的 block 數)，因此跨 block 同步/合作不會 deadlock。用 cudaLaunchCooperativeKernel() 啟動(跨裝置用 cudaLaunchCooperativeKernelMultiDevice())，配合 grid.sync()。不可 oversubscribe:一般 grid 的 block 採 transparent scalability 分波排程、順序不保證，若 block 數超過 GPU 容量，做 grid-wide barrier 時等待中的 block 永遠不會上場 → deadlock；cooperative kernel 用「犧牲可任意擴張的 block 數」換「全 grid 同時駐留」。

早期 GPU 的 FP64 約比 FP32 慢 8 倍(被 HPC 社群詬病)。Fermi 及後繼把 FP64 運算單元強化到約 FP32 的 一半速度(2x slower)，讓移植 CPU 數值程式的開發者不必再費力評估能否降成 single-precision，大幅降低移植成本。half-precision (FP16) 硬體支援在 Pascal 引入，Ampere (A100) 再透過 tensor cores 把 FP16 吞吐推到極致。

(a) $156/19.5=8$ → FP16 約為 FP32 的 8x 吞吐。(b) 還在 記憶體頻寬 受益:16-bit 資料的搬運量只有 32-bit 的一半,對 bandwidth-bound 的 kernel 直接減半頻寬壓力(也省能耗)。前提是資料本身能容忍該表示範圍與 round-off(先確認數值穩定性再降精度)。

(a) 兩處:malloc/free → cudaMallocManaged()/cudaFree()；qsort() → launch 一個平行 sort kernel + cudaDeviceSynchronize()(仍須有 parallel qsort kernel,但 host code 改動極小)。(b) 需 Pascal 起的 page fault handling + system-wide VA:因為 host 與 device 用同一指標值指同一變數,host 建立的指標鏈 device 也能跟著走(反之亦然),且資料不必放在 zero-copy memory;page fault 會按需把 page migrate 到 GPU。這對 CAD 等需數百 GB「in core」資料的應用特別關鍵。

舊系統一次只能跑一個 grid,且 grid 要夠大才能填滿 device → 兩難:等 remote work 才動作會讓 device 閒置、低利用率;急著先做 local work 又會讓 remote work 卡在大 grid 後面、latency 變高。Fermi 起的 simultaneous multiple grids(配合 Kepler 的多 hardware queues)解決:用較小的 grid size 提交工作,並可分 priority;當高優先的 remote work 到達時,可低延遲立即插入執行,而不必卡在一大塊 local computation 之後。

存取方式:zero-copy 永遠走 PCIe 直接讀、無遷移(頻寬 <10% global memory,故只能稀疏存取);unified memory 會按需 migrate page 到存取端,或 map-on-access。遍歷指標鏈:zero-copy 只在「全部記憶體都用 cudaHostAlloc() 配置」時才能 traverse;unified memory(Pascal)可遍歷任意 host 建立的 linked structure,因 host/device 共用同一指標值。資料量上限:zero-copy 受限於 host pinned memory,unified memory 涵蓋全系統(CPU + 多 GPU)的 VA。Coherence:zero-copy 無(手動),unified memory 由硬體 + page fault 自動維護。因此 unified memory 不再受「能 copy/pin 到 device 的資料量」限制,還能自動把熱資料搬到存取端,是更通用的機制。

PC sampling(CUDA 7.5)提供指令級 profiling,告訴你哪一行程式碼最耗時(熱點);但它不知道該 kernel 是否與 CPU 活動重疊。Critical path analysis(CUDA 8 Visual Profiler)則找出哪一個 kernel/API 真正決定整體執行時間,非關鍵路徑活動灰階淡化。為何最長 kernel 未必該優化:如 Fig 22.2,Kernel X 雖長卻完全與 CPU 活動 A 重疊,只加速 X 而不同時加速 A → speedup ≈ 1;反觀較短的 Kernel Y 在 critical path 上、CPU 正空等它,加速 Y 才真能縮短整體時間。這正是 Amdahl 思維:只有優化「真正占用串列關鍵路徑時間」的部分才有 speedup,故兩工具互補,缺一不可。

早期 GPU 的 atomic 慢且功能受限,開發者常被迫做演算法轉換(prefix-sum / sort)來避開 atomic,但這些轉換會增加 kernel 啟動次數與總運算量,有時甚至要把 collective/多 block 更新工作回丟 host CPU(增加 CPU↔GPU 傳輸)。隨 Fermi→Kepler→Maxwell atomic 逐代變快、更通用(Maxwell 再強化 shared-memory atomic 吞吐),直接 atomic scatter 變得可接受,減少了改用 scan/sort 或回丟 CPU 的需要。但「快的 atomic」≠ 不用優化:高競爭 (contention) 仍會嚴重拖慢效能,privatization、coarsening、aggregation 等技巧依然有效(見 09-Parallel-Histogram/02-Histogram-Optimizations-Privatization-Coarsening-Aggregation)。

記憶體模型演進時間軸

Kernel 執行控制能力(首見世代)

吞吐量 / 頻寬硬體進步:DP 8x→1/2(Fermi)；FP16 tensor 156 vs 19.5 TFLOPS(A100)；predication(Fermi)；可配置 cache/shared(Fermi)；atomic Fermi→Kepler→Maxwell 漸快；HBM2 ≤3x、NVLink ≤5x(Pascal)。

生產力工具:unified device memory space(Fermi,把晶片內 g/l/s 位址統一,勿與 managed memory 混淆)；OpenACC / Thrust / CUDA FORTRAN / C++AMP；PC sampling(CUDA 7.5,哪一行) + critical path analysis(CUDA 8,哪一個 kernel)。

一句話總結:Ch.22 的主線是「讓指標愈來愈通用、讓 GPU code 愈來愈像 CPU code」——記憶體模型從手動 cudaMemcpy 走向單一指標自動遷移,kernel 能力從強制 inline 走向 composable,硬體吞吐拉低「easy performance」門檻,工具從「最熱的行」走向「真正在 critical path 上的 kernel」。