Unified 與 System Memory (Unified and System Memory)

重點總覽

異質系統有多個實體記憶體（host DRAM、每個 GPU 各自的 DRAM）。資料位於存取它的 processor 記憶體中時效能最佳。CUDA 提供明確管理記憶體放置的 API，但繁瑣；本筆記介紹簡化配置、放置與遷移的特性。

Unified Virtual Address Space（UVA）

單一 OS process 內，所有 host memory 與所有 GPU 上的 global memory 共用一個虛擬位址空間。無論用 CUDA API（cudaMalloc、cudaMallocHost）或 system API（new、malloc、mmap）配置，記憶體都落在這個位址空間中，CPU 與每個 GPU 各佔一段獨特範圍。

• 任何記憶體的位置（在 CPU 或哪個 GPU）都可由指標值用 cudaPointerGetAttributes() 判斷。
• cudaMemcpy*() 的 cudaMemcpyKind 可設為 cudaMemcpyDefault，由指標自動判斷複製方向。

            單一虛擬位址空間 (UVA)
 ┌───────────┬───────────┬───────────┐
 │  CPU 範圍 │ GPU0 範圍 │ GPU1 範圍 │ ...
 └───────────┴───────────┴───────────┘
   ↑ 指標值就能判斷它屬於哪個 processor

Unified Memory（managed memory）

Unified memory 是讓一種稱為 managed memory 的配置可同時被 CPU 或 GPU 程式碼存取的 CUDA 特性，且能在 CPU/GPU 間自動遷移。所有 CUDA 支援的系統都有 unified memory。

在某些系統上 managed memory 必須明確配置，方式有三：

• CUDA API cudaMallocManaged
• cudaMallocFromPoolAsync，pool 以 allocType = cudaMemAllocationTypeManaged 建立
• 帶 __managed__ specifier 的全域變數

Unified Memory Paradigms（4 種模式與判斷）

unified memory 的特性與行為隨 OS、Linux kernel 版本、GPU 硬體、GPU-CPU interconnect 而不同。可用 cudaDeviceGetAttribute 查詢三個屬性來判斷型態：

四種 paradigm：

   有 unified memory?  (所有現代 GPU 皆有 UVA + unified memory)
            │
   ConcurrentManagedAccess == 1 ?
       ├── 0 ──► Limited support（Windows/WSL/Tegra）
       └── 1 ──► Full support
                   │
            PageableMemoryAccess == 1 ?
              ├── 0 ──► 只有 CUDA 配置的 managed memory 是 unified
              └── 1 ──► 所有 system memory 都是 unified
                         │
                  UsesHostPageTables == 1 ?
                    ├── 1 ──► Hardware coherence (ATS)
                    └── 0 ──► Software coherence (HMM)

查詢與判斷邏輯的精簡程式碼：

cudaDeviceGetAttribute(&concurrentManagedAccess,
                       cudaDevAttrConcurrentManagedAccess, deviceId);
cudaDeviceGetAttribute(&pageableMemoryAccess,
                       cudaDevAttrPageableMemoryAccess, deviceId);
cudaDeviceGetAttribute(&pageableMemoryAccessUsesHostPageTables,
                       cudaDevAttrPageableMemoryAccessUsesHostPageTables, deviceId);

if (concurrentManagedAccess) {
  if (pageableMemoryAccess)
    // full support；UsesHostPageTables=1 為 hardware，0 為 software coherency
    ;
  else
    // full support，僅限 CUDA 配置的 managed allocations
    ;
} else {
  // limited support: Windows, WSL, or Tegra
}

此即把上面的決策樹翻成程式碼，逐 GPU（cudaGetDeviceCount/cudaSetDevice）查詢能力。

Full Unified Memory Feature Support

大多數 Linux 系統有 full unified memory 支援。

• 若 cudaDevAttrPageableMemoryAccess = 1：所有 system memory（無論 CUDA API 或 system API 配置）都以 full feature 的 unified memory 運作，包含 mmap 建立的 file-backed 配置。
• 若 = 0：只有 CUDA 配置為 managed 的記憶體才是 unified；system API 配置的記憶體未被管理，不保證能從 GPU kernel 存取。

full support 下 managed 配置的一般行為：

• 通常配置在**第一次被觸碰（first touch）**的 processor 記憶體空間。
• 當被目前所在以外的 processor 使用時，通常會遷移。
• 遷移／存取粒度：memory page（software coherence） 或 cache line（hardware coherence）。
• 允許 oversubscription：可配置超過 GPU 實體可用量的 managed memory。

Hardware Coherency（ATS）

在如 Grace Hopper、Grace Blackwell 這類使用 NVIDIA CPU 且 CPU-GPU 以 NVLink Chip-to-Chip（C2C） 相連的硬體上，可使用 Address Translation Services（ATS）。此時 cudaDevAttrPageableMemoryAccessUsesHostPageTables = 1。

除了對所有 host 配置都有 full unified memory 支援外，ATS 還提供：

• 駐留在 GPU 的 managed 配置（如 cudaMallocManaged）可從 CPU 不經遷移直接存取（cudaDevAttrDirectManagedMemAccessFromHost = 1）。
• CPU-GPU 連結支援 native atomics（cudaDevAttrHostNativeAtomicSupported = 1）。
• 硬體 coherence 相較 software coherence 可提升效能。

HMM（Software Coherency）

Heterogeneous Memory Management（HMM） 是 Linux（需適當 kernel 版本）的功能，提供 software-coherent 的 full unified memory，把 ATS 的部分便利帶給 PCIe-connected GPU。

• 需 Linux Kernel 6.1.24、6.2.11、或 6.3 以上。
• 查詢 addressing mode：

$ nvidia-smi -q | grep Addressing
Addressing Mode : HMM

• HMM 可用時，full unified memory 受支援且所有 system 配置都隱含是 unified memory。
• 若系統同時有 ATS，HMM 停用、改用 ATS。

Limited Unified Memory Support

在 Windows（含 WSL） 與部分 Tegra 系統上，只有 unified memory 的有限子集。managed memory 仍可用，但 CPU/GPU 間遷移行為不同：

• managed memory 先配置在 CPU 實體記憶體。
• 以大於虛擬記憶體 page 的粒度遷移。
• 在 GPU 開始執行時遷移到 GPU。
• GPU 活動期間 CPU 不可存取 managed memory。
• 在 GPU 同步時遷回 CPU。
• 不允許 oversubscription。
• 只有 CUDA 明確配置為 managed 的記憶體才是 unified。

Memory Advise 與 Prefetch

程式設計者可給 NVIDIA Driver 提示，協助最大化效能。

Page-Locked Host Memory（pinned memory）

cudaMallocHost 配置的是 host 上的 page-locked memory（又稱 pinned memory）。用 malloc、new、mmap 等傳統機制配置的 host 記憶體不是 page-locked，可能被 OS 換出（swap）到磁碟或實體搬移。

• async copy（CPU↔GPU）必須使用 page-locked memory。
• 也能提升 synchronous copy 的效能。
• 可 map 給 GPU 供 kernel 直接存取。

Mapped Memory（zero-copy）

在有 HMM 或 ATS 的系統上，所有 host memory 都能用 host 指標直接從 GPU 存取。當 ATS/HMM 不可用時，可把 host 配置 map 進 GPU 記憶體空間讓 GPU 存取。Mapped memory 一律是 page-locked。

範例 kernel（直接操作 mapped host memory）：

__global__ void copyKernel(float* a, float* b)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    a[idx] = b[idx];
}

  cudaMallocHost / cudaHostAlloc：用同一個 host 指標
    GPU kernel ──(PCIe / NVLink C2C)──► [ Host (pinned) memory ]
                                          ↑ 資料始終留在 CPU 端

  cudaHostRegister：需轉換成 device 指標
    malloc 配置 ──► cudaHostRegister ──► cudaHostGetDevicePointer
    kernel 用 devPtr ─(interconnect)─► host memory

cudaMallocHost / cudaHostAlloc

用 cudaMallocHost 或 cudaHostAlloc 配置的 host 記憶體會自動被 mapped。回傳的指標可直接用在 kernel 中存取 host 上的記憶體（經 CPU-GPU interconnect）。

// cudaMallocHost：指標可直接傳給 kernel
CUDA_CHECK(cudaMallocHost(&a, vLen*sizeof(float)));
CUDA_CHECK(cudaMallocHost(&b, vLen*sizeof(float)));
int blocks = vLen / threads;
copyKernel<<<blocks, threads>>>(a, b);   // 直接用 host 指標

// cudaHostAlloc：用 flag 指定 mapped
CUDA_CHECK(cudaHostAlloc(&a, vLen*sizeof(float), cudaHostAllocMapped));
copyKernel<<<blocks, threads>>>(a, b);

重點：cudaHostAlloc 需傳 cudaHostAllocMapped 之類的 flag；兩者回傳的 host 指標都能直接在 kernel 使用。

cudaHostRegister

當 ATS/HMM 不可用時，system allocator 配置的記憶體仍可用 cudaHostRegister map 供 kernel 直接存取。但與 CUDA API 配置不同，不能用 host 指標在 kernel 存取，必須用 cudaHostGetDevicePointer() 取得 device 區域的指標，並在 kernel 中使用該指標。

a = (float*)malloc(vLen*sizeof(float));
CUDA_CHECK(cudaHostRegister(a, vLen*sizeof(float), 0));
CUDA_CHECK(cudaHostGetDevicePointer((void**)&devA, (void*)a, 0));
// kernel 必須使用 devA（device 指標），不可用 a
copyKernel<<<blocks, threads>>>(devA, devB);

重點：cudaHostRegister + cudaHostGetDevicePointer 是用 host 指標 a 與 device 指標 devA 兩個不同指標。

Unified Memory vs Mapped Memory

考試/測驗重點

Related Notes

• 02-Programming-GPUs/02-CUDA-Cpp-Memory-Management
• 02-Programming-GPUs/03-CUDA-Cpp-Sync-and-Workflow
• 02-Programming-GPUs/07-SIMT-Device-Memory-Spaces
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/15-Async-Callbacks-Ordering-Graphs
• 02-Programming-GPUs/17-NVCC-Compiler
• 01-Introduction-to-CUDA/04-GPU-Memory-Hierarchy
• 01-Introduction-to-CUDA/05-CUDA-Platform
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps