Virtual Memory Management

重點總覽

動機與設計（為何需要 VMM）

CUDA 一般配置（如 cudaMalloc）直接回傳可用的 GPU 位址，位址與實體記憶體綁死。VMM 仿照 OS 的虛擬記憶體模型，分兩階段：先保留一段連續虛擬位址（不配實體），用到時再提交/映射實體儲存。

• Fine-grained 控制：非連續的實體 chunk 可映射成連續 VA，降低碎裂、提升大型工作負載（如 DNN 訓練）利用率。
• 配置/釋放高效：把 VA 保留與實體配置解耦，可先 reserve 大區、再依需求 map，避免昂貴的 copy/realloc。
• 動態成長：類似 CPU 的 realloc 或 std::vector，不需搬移全部資料即可擴張。
• 分散式多 GPU：解耦 VA 與實體後可建立統一虛擬位址空間，資料動態映射到不同 GPU；NCCL、NVSHMEM 即建構於此。

Preliminaries：定義與支援查詢

關鍵名詞

• CU_MEM_HANDLE_TYPE_FABRIC（CUDA 12.4+）：在支援平台且 IMEX daemon 執行時，可用 MPI 等任意機制做 intra-node 與 inter-node 共享，讓多節點 NVLink 系統的 GPU 互相映射記憶體。

逐 device 查詢支援

int deviceSupportsVmm;
cuDeviceGetAttribute(&deviceSupportsVmm,
    CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, device);
// fabric：CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_FABRIC_SUPPORTED
// multicast：CU_DEVICE_ATTRIBUTE_MULTICAST_SUPPORTED

• 功能可用性依「GPU 架構、driver 版本、軟體庫」而異，使用前務必逐 device 查詢。
• 除了 handle type 用 CU_MEM_HANDLE_TYPE_FABRIC 且不需 OS 原生 IPC 機制外，fabric memory 的用法與其他 handle type 無異。

IMEX channels 環境檢查（fabric 共享）

# 1) 確認 character device 存在（取得 major number，如 234）
cat /proc/devices | grep nvidia    # 出現 nvidia-caps-imex-channels
# 2) exporter 與 importer 須能存取同一 channel 檔
mknod /dev/nvidia-caps-imex-channels/channelN c <major_number> 0

• channel 檔（如 /dev/nvidia-caps-imex-channels/channel0）代表個別 IMEX channel，須由系統管理員建立。
• 若指定 NVreg_CreateImexChannel0 模組參數，driver 預設可自建 channel0。

API Overview 與工作流程

VMM 是非常低階的 API，必須直接用 CUDA Driver API，可用於單節點與多節點。使用者需具備 OS 虛擬記憶體、記憶體階層、IPC 方法（socket/message passing）、與記憶體存取權安全的基本知識。

工作流程的核心是「在 source device 配置實體記憶體後，用 handle 把必要資訊傳給 target」。export 的 handle 可為 OS-specific（限單節點 IPC）或 fabric-specific（單/多節點，需 IMEX）。share 用何種 IPC 由開發者自選；之後雙方各自 reserve VA、map、設 access。

              ┌─ 評估環境設定 ─┐
              │  選擇 handle 類型 │
        ┌─────┴─────┐     ┌─────┴─────┐
   OS-specific handle   Fabric handle
   （單節點 IPC）        （單/多節點，需 IMEX）
        └─────┬─────┘     └─────┬─────┘
              └────────┬────────┘
        ─── 後段操作對兩種 handle 完全相同 ───
            Reserve VA → Map → 設定 Access Rights

Unicast Memory Sharing：五步驟

Allocate & Export → Share & Import → Reserve & Map → Access Rights → Releasing。
此流程同時涵蓋「VA → GPU 實體位址（PA）→ 網路 Fabric 位址（FA）」的映射。

1. Allocate & Export

cuMemCreate 配置實體記憶體 backing。回傳的 CUmemGenericAllocationHandle 不是指標、尚不可存取。CUmemAllocationProp 描述位置、是否分享、實體屬性；size 須對齊以 cuMemGetAllocationGranularity 查到的 granularity。

CUmemAllocationProp prop = {};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
prop.location.id = device;
prop.requestedHandleType = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR; // 或 _FABRIC
size_t granularity = 0;
cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM);
size_t padded_size = ROUND_UP(size, granularity);   // 對齊到 granularity
CUmemGenericAllocationHandle allocHandle;
cuMemCreate(&allocHandle, padded_size, &prop, 0);

• OS-specific 與 fabric 版本唯一差異是 requestedHandleType（..._POSIX_FILE_DESCRIPTOR vs ..._FABRIC）。
• 配置屬性可用 cuMemGetAllocationPropertiesFromHandle 查詢。

export 把 handle 轉成可分享形式：

int fd;
cuMemExportToShareableHandle(&fd, handle, CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR, 0);
// fabric：
CUmemFabricHandle fh;
cuMemExportToShareableHandle(&fh, handle, CU_MEM_HANDLE_TYPE_FABRIC, 0);

• OS-specific handle 要求所有 process 屬同一 OS；fabric handle 要求 sysadmin 啟用 IMEX channels。
• 範例：memMapIpcDrv。

2. Share & Import

export 後須用 IPC 把 handle 傳給接收方，方法由開發者自選。OS-specific IPC（如 Unix domain socket 配 SCM_RIGHTS、Win32 DuplicateHandle）效能高但限同機、不可攜；fabric IPC（如 MPI_Send/MPI_Recv 傳 CUmemFabricHandle）簡單可攜但需系統層支援。

// Linux socket：以 SCM_RIGHTS 經 cmsg 傳送 fd（sendmsg / recvmsg）
// Fabric：MPI_Send(&fh, sizeof(CUmemFabricHandle), MPI_BYTE, 1, 0, MPI_COMM_WORLD);
cuMemImportFromShareableHandle(handle, (void*)&fd, CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR);
// fabric：cuMemImportFromShareableHandle(handle, (void*)&fh, CU_MEM_HANDLE_TYPE_FABRIC);

• import 在接收端用 handle 重建必要的記憶體物件，得到一個可後續 map 的 CUmemGenericAllocationHandle。

3. Reserve & Map

VMM 中「位址」與「記憶體」分離。先用 cuMemAddressReserve 保留一段 VA（至少等於要放入的所有實體配置 size 總和），再用 cuMemMap 把 VA 與實體 handle 綁定。

CUdeviceptr ptr;
cuMemAddressReserve(&ptr, size, 0, 0, 0);   // alignment=0 用預設
// allocHandle 來自 cuMemCreate 或 cuMemImportFromShareableHandle
cuMemMap(ptr, size, 0, allocHandle, 0);

• cuMemAddressReserve / cuMemAddressFree 概念上等同 Linux 的 mmap/munmap、Windows 的 VirtualAlloc/VirtualFree。
• 可把不同 device 的配置擺進同一段連續 VA（只要保留夠大）；用 cuMemUnmap 解耦。
• 可對同一 VA 反覆 map/unmap，但不得在已映射的 VA 上重複建立映射。cuMemAddressFree 前須先全部 unmap。

VMM 位址 vs 記憶體分離：

  Reserve:   [────── VA range (無實體) ──────]
  Map:       [─VA─]──cuMemMap──►[ allocHandle(實體) ]
  Access:    設 cuMemSetAccess 後才可被 kernel 讀寫

4. Access Rights

cuMemMap 不會讓位址可存取；未設 access 就被 kernel 存取會 crash。必須在 source 與存取端 device 上呼叫 cuMemSetAccess。

CUmemAccessDesc accessDesc = {};
accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
accessDesc.location.id   = device;
accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
cuMemSetAccess(ptr, size, &accessDesc, 1);   // 讓位址可存取

• 相對 cudaEnablePeerAccess 強制映射所有配置，VMM 的 allocation 粒度 access 讓 peer 映射開銷極小。
• 範例：vectorAddMMAP。

5. Releasing

source 與 target 都必須依序呼叫三個函式，確保實體記憶體與 VA 都乾淨釋放。

cuMemUnmap(ptr, size);        // 1) 解除 VA 與實體的映射
cuMemRelease(handle);         // 2) 釋放實體記憶體還給系統
cuMemAddressFree(ptr, size);  // 3) 釋放 VA 範圍供未來重用

Multicast Memory Sharing

Multicast Object 結合 VMM，使支援的 NVLink+NVSwitch GPU 能用 NVLink SHARP 做 in-fabric 運算（broadcast、reduction）。N 個 GPU 組成 multicast team，每個 GPU 各自用實體記憶體 back 一份 replica，因此一個 multicast object 有 N 份本地 replica。multimem PTX 指令對所有 replica 生效。

Multicast team（N GPUs，NVSwitch 連接）：

  mcHandle ──┬─ GPU0: replica0（本地實體）
             ├─ GPU1: replica1（本地實體）
             └─ GPUn: replican（本地實體）
  multimem PTX 一次操作 → 同時作用於所有 replica

四步驟：

CUmemAllocationProp mcProp = {};
mcProp.numDevices  = numDevices;
mcProp.handleTypes = CU_MEM_HANDLE_TYPE_FABRIC; // 單節點可用 POSIX_FILE_DESCRIPTOR
cuMulticastGetGranularity(&granularity, &mcProp, CU_MEM_ALLOC_GRANULARITY_MINIMUM);
mcProp.size = ROUND_UP(size, granularity);
cuMulticastCreate(&mcHandle, &mcProp);

綁定後再像 unicast 一樣 reserve VA、map mcHandle、設 access，並可同時建立 unicast 與 multicast 映射到同一實體（一致性見 Virtual Aliasing）。最後在 kernel 用 inline PTX 的 multimem 指令：

// 對所有 replica 做 atomic add reduction（需 CC 9.0+）
cuda::ptx::multimem_red(cuda::ptx::release_t, cuda::ptx::scope_sys_t,
                        cuda::ptx::op_add_t, arrival_counter_mc, n);
cuda::ptx::fence_proxy_alias();   // mc 與 uc 存取同一記憶體間需 fence
asm volatile("multimem.ld_reduce.relaxed.sys.global.add.f32 %0, [%1];"
             : "=f"(l2_norm_sum) : "l"(partial_l2_norm_mc) : "memory");

• multicast 與 unicast 映射到同一記憶體時，兩者間須加 fence.proxy.alias（不同 proxy 的存取需排序）。
• multimem.ld_reduce 從所有 replica 做 atomic load reduction；它不提供排序故可用 relaxed。
• 完整範例：Multi GPU Programming Models repo 的 multi_node_p2p，示範 NVSHMEM 在 NVLink domain 內部如何運作。

Advanced Configuration

Memory Type 與 Compressible Memory

cuMemCreate 可透過 CUmemAllocationProp::allocFlags opt-in 特殊記憶體特性（須 device 支援）。Compressible memory 可加速非結構稀疏等可壓縮資料的存取，省下 DRAM 頻寬、L2 讀頻寬與 L2 容量。

int compressionSupported = 0;
cuDeviceGetAttribute(&compressionSupported,
    CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED, device);
prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC;  // 配置時 opt-in

Virtual Aliasing Support

對同一份實體配置用多次 cuMemMap（不同 VA）建立多個「proxy」即 virtual aliasing。除非 PTX ISA 另有說明，對一個 proxy 的寫入，在寫入端操作（grid launch、memcpy、memset…）完成前，對其他 proxy 視為不一致、不一致快取。

__global__ void foo(char *A, char *B) {  // A、B 為同一記憶體的 alias
    *A = 0x1;
    printf("%d\n", *B);   // 未定義行為！可能讀到舊值或中間值
}

• 合法寫法一：用 stream/event 讓兩個 kernel 單調排序，讀的 kernel（與 memcpy）等寫的 foo1 完成後才執行。
• 合法寫法二：同一 kernel 內若必須跨 proxy 存取，在兩次存取間插入 fence.proxy.alias。

__global__ void foo(char *A, char *B) {
    *A = 0x1;
    cuda::ptx::fence_proxy_alias();   // 建立 alias proxy 間的排序
    printf("%d\n", *B);               // *B == *A == 0x1
}

OS-Specific Handle Details for IPC

cuMemCreate 時可用 CUmemAllocationProp::requestedHandleTypes 預先標記某配置供 IPC 或 graphics interop。

#if defined(__linux__)
  cuDeviceGetAttribute(&ok, CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR_SUPPORTED, device);
  prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;
#else
  cuDeviceGetAttribute(&ok, CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_HANDLE_SUPPORTED, device);
  prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_WIN32;
  prop.win32HandleMetaData  = /* LPSECURITYATTRIBUTES */;  // 定義可被轉移到哪些 process
#endif

• Windows 設 CU_MEM_HANDLE_TYPE_WIN32 時，必須在 win32HandleMetaData 提供 LPSECURITYATTRIBUTES，界定 export 配置可轉移的範圍。
• export 前務必先查詢該 handle type 是否被 device 支援。

考試/測驗重點

Related Notes

• 04-CUDA-Features/19-Interprocess-Communication
• 04-CUDA-Features/21-Extended-GPU-Memory
• 04-CUDA-Features/06-Stream-Ordered-Memory-Allocator
• 04-CUDA-Features/17-L2-Cache-Control
• 04-CUDA-Features/18-Memory-Synchronization-Domains
• 03-Advanced-CUDA/08-CUDA-Driver-API
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps