Unified Memory：完整支援深入 (Unified Memory Full Support)

重點總覽

Full CUDA Unified Memory Support 概觀

統一記憶體共有四種範式，本筆記聚焦前三種「完整支援」：

• 完整支援明確的 managed allocation（cudaMallocManaged）
• 完整支援所有 allocation + software coherence
• 完整支援所有 allocation + hardware coherence

前三者的行為與程式模型「非常相似」，差異會逐點標註；第四種「有限支援」（Windows、WSL、Tegra）另行討論，見 04-CUDA-Features/02-Unified-Memory-Platforms-and-Hints。

In-Depth Examples：存取各種 host 記憶體

具備完整支援的系統，允許 device 存取 host process 擁有的「任何」記憶體。以下 kernel 印出輸入字元陣列前 8 個字元，可用多種 host 記憶體來源呼叫：

__global__ void kernel(const char* type, const char* data) {
  static const int n_char = 8;
  printf("%s - first %d characters: '", type, n_char);
  for (int i = 0; i < n_char; ++i) printf("%c", data[i]);
  printf("'\n");
}

可傳入的記憶體來源涵蓋：malloc/new heap、cudaMallocManaged、stack 區域變數、file-scope static、global-scope 變數、以及由第三方函式庫擁有的 extern 變數（該函式庫可完全不接觸 CUDA）。

int global_variable;                 // global scope, __host__-only
__global__ void kernel(int* p) { printf("%d\n", *p); }
int main() {
  // 在 pageableMemoryAccess=1 系統上，傳「位址」而非直接存取變數
  kernel<<<1, 1>>>(&global_variable);
}

重點：直接在 device code 中讀 global_variable 會編譯失敗；必須改傳其位址。

File-backed Unified Memory

由於 device 可存取 host process 的任意記憶體，因此能直接存取 file-backed 記憶體。做法是用 mmap 把檔案映射進記憶體，再把回傳指標直接交給 kernel：

int fd = open(INPUT_FILE_NAME, O_RDONLY);
char* mapped = (char*)mmap(0, file_stat.st_size, PROT_READ, MAP_PRIVATE, fd, 0);
kernel<<<1, 1>>>("file-backed", mapped);   // GPU 直接讀檔案映射記憶體
cudaDeviceSynchronize();
munmap(mapped, file_stat.st_size);

實體檔案與 memory-backed 檔案皆適用。

IPC with Unified Memory

許多應用偏好一個 process 管一張 GPU，但仍需用統一記憶體（如 over-subscription）並從多張 GPU 存取。

• CUDA IPC「不支援」managed memory：這類記憶體的 handle 無法透過任何 IPC 機制分享。
• 完整支援系統上，system-allocated memory 具 IPC 能力；一旦分享出去，程式模型與 File-backed Unified Memory 相同。
• Linux 下可建立 IPC-capable system-allocated memory 的方式：mmap 搭 MAP_SHARED、POSIX IPC API、memfd_create。

詳見 04-CUDA-Features/19-Interprocess-Communication。

Memory Paging 與 Page Sizes

所有支援統一記憶體的系統都使用 virtual address space 與 memory paging，因此存在兩類頁：

• Virtual page：每個 process 由 OS 追蹤的固定大小虛擬連續區塊，可映射到實體記憶體；page 與映射綁定（同一虛擬位址可用不同 page size 映射）。
• Physical page：處理器 MMU 支援的固定大小實體連續區塊，virtual page 映射至此。

映射關係由 page table 追蹤，每筆映射稱 PTE；處理器以 TLB 快取 page table 以加速 virtual→physical 位址轉換。效能調校的兩大關鍵：選擇 virtual page size、以及系統是「合併 page table」還是「各自獨立」。

Choosing the Right Page Size

• GPU 的 TLB miss「明顯比 CPU 昂貴」：若 GPU thread 頻繁隨機存取以「小頁」映射的統一記憶體，會比用「夠大頁」映射時慢很多。
• CPU 隨機存取小頁大區域也會變慢，但程度較輕；應用可用此「較小的減速」換取「較少的碎片」。

CPU/GPU Page Tables：Hardware vs Software Coherency

Hardware-coherent (Grace Hopper)        Software-coherent (HMM 等)
  CPU ─┐                                  CPU ── 自己的 page table
       ├─ 邏輯合併 page table             GPU ── 自己的 page table
  GPU ─┘                                   存取對方實體頁 → page fault + migration
  cache-line 粒度一致                      page 粒度一致

• Hardware-coherent：GPU 直接使用 CPU 為該記憶體建立的 PTE。若該 PTE 用 CPU 預設 4KiB/64KiB 小頁，存取大區域會引發大量 TLB miss 與明顯減速。
• Software-coherent：CPU/GPU 各有獨立 page table，以 page fault 保證一致性。一次 fault 需：(1) 讓原擁有者不能再存取（刪除/更新 PTE）；(2) 讓請求者可存取（建立/更新 PTE 使其有效）；(3) 把實體頁 migration 到請求者——成本與 page size 成正比。

Mixing Hardware and Software Coherency

部分 hardware-coherent 系統（如 NVIDIA DGX Station）也支援安裝獨立的非一致性 GPU。兩類 GPU 共享同一統一位址空間，但行為不同：

• hardware-coherent GPU 的存取沿用硬體一致性，較少 fault、盡量遠端映射。
• 軟體一致性的 discrete GPU 存取會有「更多 page fault 與 migration」。
• cudaHostRegister 等 host 記憶體 API 在 mixed 系統上對 software-coherent GPU 改用「軟體鏡射 CPU page table」而非 pinned 映射；在記憶體壓力下可能罕見地多出 page fault。

Direct Unified Memory Access from the Host

部分裝置具備從 host 對「GPU-resident 統一記憶體」做 coherent 讀/寫/atomic 的硬體支援，屬性 cudaDevAttrDirectManagedMemAccessFromHost=1（所有 hardware-coherent 系統對 NVLink 連接的 device 皆設此屬性）。此時 host 直接存取 GPU 記憶體「無 page fault、無 migration」。

int *ret;
cudaMallocManaged(&ret, 1000 * sizeof(int));
cudaMemLocation location = {.type = cudaMemLocationTypeHost};
// managed memory 需此 hint 才能啟用「無 fault」的 direct host access
cudaMemAdvise(ret, 1000 * sizeof(int), cudaMemAdviseSetAccessedBy, location);
write<<<1, 1000>>>(ret, 10, 100);   // 資料 populate 於 GPU 記憶體
cudaDeviceSynchronize();
for (int i = 0; i < 1000; i++) printf("%d: A+B = %d\n", i, ret[i]);  // CPU 讀
append<<<1, 1000>>>(ret, 10, 100);  // GPU 再次存取

兩種系統行為對比：

Host Native Atomics

部分裝置（含 hardware-coherent 系統中 NVLink 連接的 device）支援「對 CPU-resident 記憶體的硬體加速 atomic」，意即 host 記憶體的 atomic「不需用 page fault 模擬」。這類裝置屬性 cudaDevAttrHostNativeAtomicSupported=1。

Atomic Accesses 與 Synchronization Primitives

CUDA 統一記憶體支援 host 與 device thread「所有」可用的 atomic 操作，讓所有 thread 能並行存取同一共享位址協作。libcu++ 提供多種為 host/device 並行調校的異質同步原語：cuda::atomic、cuda::atomic_ref、cuda::barrier、cuda::semaphore 等。

__global__ void kernel(int* ptr) { cuda::atomic_ref{*ptr}.store(2); }
// host 與 device 並行對同一 file-backed ptr 做 atomic（僅 hardware-coherent 系統合法）
kernel<<<1, 1>>>(ptr);
while atomic_ref{*ptr}.load() != 2;   // 最終必為 2

• software-coherent 系統上，對統一記憶體的 atomic 可能引發 page fault、帶來顯著 latency；但並非「所有」GPU 對 CPU 記憶體的 atomic 都如此——nvidia-smi -q | grep "Atomic Caps Outbound" 列出的操作可避免 fault。
• hardware-coherent 系統上，host/device 間 atomic 不需 page fault，但仍可能因「任何記憶體存取都可能 fault 的其他原因」而 fault。

詳見 04-CUDA-Features/11-Asynchronous-Barriers-Deep-Dive 的 barrier 與 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy 的 atomic 入門。

其他效能調校重點

Memcpy() / Memset() 行為

cudaMemcpy*() 與 cudaMemset*() 接受任何統一記憶體指標。對 cudaMemcpy*() 而言，cudaMemcpyKind 方向是「效能 hint」，當任一引數為統一記憶體指標時影響更大。建議：

• 已知統一記憶體實體位置時，使用「準確」的 cudaMemcpyKind。
• 不確定方向時，cudaMemcpyDefault 優於「不準確」的 hint。
• 一律對「已 populate（已初始化）」的 buffer 操作；勿用這些 API 來初始化記憶體。
• 若兩個指標都指向 system-allocated 記憶體，避免 cudaMemcpy*()，改以 kernel 或 std::memcpy。

統一記憶體 allocator 概觀

Access Counter Migration（僅 hardware-coherent）

access counters 追蹤 GPU 對「他處理器記憶體」的存取頻率，把頁搬到存取最頻繁者的實體記憶體（可在 CPU↔GPU、peer GPU↔GPU 之間 guide migration）。

• CUDA 12.4 起 system-allocated memory 支援 access counters；file-backed 記憶體「不」依存取 migration。
• 用 cudaMemAdviseSetAccessedBy（對應 device id）開啟；開啟後可用 cudaMemAdviseSetPreferredLocation 設為 host 來防止 migration。
• cudaMallocManaged 預設採 fault-and-migrate 機制。

避免 CPU 頻繁寫入 GPU-resident 記憶體

許多 CPU 架構所有記憶體操作（含 write）都須經 cache。若記憶體常駐於 GPU，CPU 頻繁寫入會 cache miss，先把資料 GPU→CPU 搬回再寫入。建議「寫入 CPU-resident 記憶體、由 device 直接讀」：用 cudaMemAdviseSetPreferredLocation + cudaMemAdviseSetAccessedBy（皆 cudaMemLocationTypeHost）把資料釘在 host。

善用對 system memory 的非同步存取

device 與 host 交換結果的三種方式：

1. device 寫 GPU 記憶體 → cudaMemcpyAsync 傳到 host → host 讀。
2. device 直接寫 CPU-resident 記憶體 → host 讀。
3. device 寫 GPU 記憶體 → host 直接存取。

考試/測驗重點

Related Notes

• 04-CUDA-Features/02-Unified-Memory-Platforms-and-Hints
• 04-CUDA-Features/06-Stream-Ordered-Memory-Allocator
• 04-CUDA-Features/11-Asynchronous-Barriers-Deep-Dive
• 04-CUDA-Features/19-Interprocess-Communication
• 04-CUDA-Features/20-Virtual-Memory-Management
• 04-CUDA-Features/21-Extended-GPU-Memory
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps