Host/Device 互動模型與記憶體演進 (Host/Device Interaction Model)

重點總覽 (Overview)

本節追蹤 CPU-GPU 記憶體模型 從「Host/Device 各自獨立、必須 cudaMemcpy() 搬資料」演進到「單一指標、自動遷移」的歷史。每一步都是為了解決早期簡單模型的兩大痛點:I/O 裝置只認 host memory 與 大型資料結構塞不進 device memory。

Host/Device 簡單模型的限制 (Limitations of the Simple Model)

前面章節 (見 02-Heterogeneous-Data-Parallel-Computing/02-Vector-Addition-Host-Code-Device-Memory-and-Data-Transfer) 假設的模型:device memory (CUDA global memory) 與 host (system) memory 完全分離,必須用 cudaMemcpy() 來回搬資料。

傳統模型 (simple model)
  Host Memory  --cudaMemcpy(H2D)-->  Device Memory
   (system)    <--cudaMemcpy(D2H)--   (GPU global)
       |                                   |
   I/O 裝置                              Kernel
 (disk/NIC 只認 host memory)

兩大痛點:

• I/O 效率:disk controller、NIC 等 I/O 裝置設計成在 host memory 上運作。資料先 H2D 再運算、運算完 D2H 才能 I/O → 增加 I/O latency、降低 throughput。
• 大型資料結構:傳統程式把大資料結構放 host memory;早期 GPU device memory 很小,逼開發者把資料 切塊 (partition) 才塞得下。
  • 例:18-Electrostatic-Potential-Map/01-DCS-Scatter-vs-Gather 的 3D 能量網格被切成 2D slices 來回搬運。
  • 有些應用根本沒有好的切法 → 最好讓 GPU 直接存取 host memory,或由 runtime 自動遷移。

Zero-Copy Memory 與 UVA (Zero-Copy Memory and Unified Virtual Addressing)

Zero-Copy Memory (CUDA 2.2, 2009)

讓 kernel 直接透過 system interconnect (PCIe) 存取 host memory,完全不呼叫 cudaMemcpy()。

// 1) 配置 pinned + mapped 的 host memory
float *h_ptr;
cudaHostAlloc((void**)&h_ptr, size, cudaHostAllocMapped);

// 2) host 指標「不能」直接傳給 kernel,要先換成 device 指標
float *d_ptr;
cudaHostGetDevicePointer((void**)&d_ptr, h_ptr, 0);

// 3) 把 device 指標傳進 kernel;kernel deref 時走 PCIe 直接讀寫 host memory
myKernel<<<grid, block>>>(d_ptr);

• Zero-copy memory 必須是 pinned host memory (見 20-Heterogeneous-Computing-Cluster/03-Overlapping-Computation-and-Communication),避免 OS 在 GPU 存取時把頁面換出。
• 同一塊實體記憶體,host 與 device 用不同指標 (h_ptr vs d_ptr) 存取。

        h_ptr (host VA)                d_ptr (device VA)
            \                            /
             \                          /
              v                        v
        ┌─────────────────────────────────┐
        │  同一塊 pinned host 實體記憶體      │  <-- 兩個不同虛擬位址
        └─────────────────────────────────┘
                 ▲ PCIe 直接存取 (慢)
              Kernel on GPU

Unified Virtual Addressing (UVA, CUDA 4, 2011)

• UVA 前:host 與 device 各有 虛擬位址空間 → 同一實體位址可能對應到 host/device 兩個不同虛擬位址 (zero-copy 即是如此)。
• UVA 後:單一共用虛擬位址空間 (UVAS),保證 每個實體位址只對應一個虛擬位址。
  • runtime 只看指標值就能判斷它指向 host 還是 device → cudaMemcpy() 不再需要指定方向 (可用 cudaMemcpyDefault)。
  • 源自 GMAC library (Gelado et al., 2010)。

大型虛擬與實體位址空間 (Large Virtual and Physical Address Spaces)

32-bit 的天花板

早期 CUDA GPU 用 32-bit 虛擬位址 + 最多 32-bit 實體位址:

• device memory 上限 4 GB;CUDA kernel 只能操作 < 4 GB 的資料集。
• 現代 CPU 用 64-bit 虛擬位址 (實際用 48 bits) → GPU 的 32-bit 虛擬位址 裝不下 CPU 位址,這也限制了 zero-copy 能支援的資料結構類型。

Kepler (2013) 起的現代虛擬記憶體

好處:

• GPU 可裝 >4 GB DRAM;kernel 可操作大型資料集。
• host 與 device 可用 完全相同的指標值 存取同一份資料 (不論在哪)。
• 跨 GPU 統一實體位址空間 → 一顆 GPU 可直接 deref 指向另一顆 GPU memory 的指標 (peer access)。

Kepler 之前 (GPU 間溝通):           Kepler 之後 (直接 peer access):
  GPU0 mem --D2D copy(host觸發)--> GPU1   GPU0 直接 deref --> GPU1 mem
    (額外記憶體 + copy overhead)            (指標傳過去即可,如 halo exchange)

統一記憶體 (Unified Memory)

Managed Memory (CUDA 6, 2013)

建立一個 CPU/GPU 共享的 managed memory pool,用 單一指標 存取。變數可駐留在 CPU 實體記憶體、GPU 實體記憶體、或兩者都有;runtime + 硬體負責 data migration 與 coherence。

移植成本極低 (Fig 22.1):CPU code 改兩處即可變 CUDA:

// CPU 版                          // Unified Memory CUDA 版
int *data = (int*)malloc(N*4);     cudaMallocManaged(&data, N*4); // 改 1
// ...fill data...                 // ...fill data... (CPU 直接寫)
qsort(data, N, 4, cmp);            sortKernel<<<g,b>>>(data, N);  // 改 2
                                   cudaDeviceSynchronize();
free(data);                        cudaFree(data);

Kepler/Maxwell 的限制 (軟體為主的實作)

• 每次 grid launch 前,CPU 改過的 managed memory 全部要 flush 到 GPU device memory。
• CPU 與 GPU 不能同時 存取同一塊 managed allocation。
• unified address space 受限於 GPU 實體記憶體大小。
• 原因:這些世代 缺乏 host/device 之間的硬體 coherence,migration 多由軟體完成。

Pascal (2016):Page Fault Handling + 49-bit VA

兩大硬體特性:

1. 49-bit 虛擬位址 → 足以涵蓋現代 CPU 的 48-bit VA + GPU 自身記憶體 → 整個系統 (所有 CPU/GPU) 成為 單一虛擬位址空間,不再受限於「能 copy 到 device 的資料量」。
2. Page fault handling:

Pascal 之前:launch 前 flush 全部 managed memory  (粗暴、慢)
Pascal 之後:on-demand,靠 page fault + page protection 做 coherence

  GPU grid 存取某 page
        |
        v
  page 不在 GPU / 已被 CPU invalidate?
        |── 是 ──> 觸發 page fault
        |            ├─ migrate page 到 GPU memory (按需),或
        |            └─ map on access 走 interconnect (偶爾存取時更快)
        └── 否 ──> 直接存取

• coherence 機制:host/device 修改 managed 變數時,用 page mapping/protection 讓對方的副本失效 (invalidate);launch 時 不必 把所有 GPU 副本更新到最新。
• system-wide:GPU/CPU 可從 CPU memory 或 其他 GPU memory fault + migrate page。
• CPU deref 指向 GPU memory 的指標也能被服務 (latency 較長) → 方便呼叫 尚未移植到 GPU 的 legacy library。

虛擬位址空間控制 (Virtual Address Space Control)

CUDA 11 提供一組 low-level API 讓程式設計師精細控制記憶體配置:

cuMemAddressReserve(&ptr, size, ...);  // 1. 保留一段虛擬位址範圍
cuMemCreate(&handle, size, &prop, 0);  // 2. 在任一 device 配置實體記憶體
cuMemMap(ptr + offset, size, 0, handle, 0); // 3. 映射到保留範圍的任一位置

• 可建立 跨多裝置的自訂資料佈局。
• 例:把一個 3D volume 配置橫跨多顆 GPU,卻仍用 單一指標 來引用整個 volume。

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 22-Advanced-Practices-And-Future-Evolution/02-Kernel-Execution-Control
• 22-Advanced-Practices-And-Future-Evolution/03-Memory-Bandwidth-and-Compute-Throughput
• 22-Advanced-Practices-And-Future-Evolution/04-Programming-Environment-and-Future-Outlook
• 02-Heterogeneous-Data-Parallel-Computing/02-Vector-Addition-Host-Code-Device-Memory-and-Data-Transfer
• 05-Memory-Architecture-And-Data-Locality/01-Memory-Access-Efficiency-and-CUDA-Memory-Types
• 20-Heterogeneous-Computing-Cluster/03-Overlapping-Computation-and-Communication