Extended GPU Memory (EGM)

#cuda #memory #extended-gpu-memory #concept

重點總覽

Extended GPU Memory（EGM）利用高頻寬的 NVLink-C2C，讓 GPU thread 能在單節點與多節點系統中高效存取「所有系統記憶體」（包含 CPU attached memory 與 HBM3）。本機存取走 NVLink-C2C；遠端存取走 GPU NVLink（部分情況也用 NVLink-C2C），並可跨 NVSwitch fabric 路由。

項目	重點
核心價值	GPU thread 以 GPU-GPU NVLink 或 NVLink-C2C 速度存取整個系統記憶體
適用平台	integrated CPU-GPU NVIDIA 系統（Arm-based CPU + C2C 互連）
三種拓撲	Single-Node Single-GPU、Single-Node Multi-GPU、Multi-Node Multi-GPU
位置識別	用 OS 指派的 NUMA node id（numaID），非 device ordinal
取得 numaID	`cuDeviceGetAttribute` + `CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID`
支援的 allocator	`cuMemCreate`（VMM）與 `cudaMemPoolCreate`（Stream Ordered）
Location type	`CU_MEM_LOCATION_TYPE_HOST_NUMA` / `cudaMemLocationTypeHostNuma`
多節點關鍵	`CU_MEM_HANDLE_TYPE_FABRIC` + IPC（export/import shareable handle）
頁面大小	EGM 以 2MB pages 映射，超大配置可能更多 TLB miss

Tip

EGM 的價值在於「路由保證走 NVLink」。把遠端 socket 的系統記憶體映射成 EGM 不但不會掉效能，存取反而更快，因為流量保證經由 NVLink 而非較慢的路徑。

概觀：EGM 是什麼

EGM 利用高頻寬 NVLink-C2C，讓 GPU 在單節點與多節點皆能高效存取所有系統記憶體。
套用於 integrated CPU-GPU NVIDIA 系統：允許「實體記憶體配置」可被 setup 內任一 GPU thread 存取。
保證所有 GPU 都能以 GPU-GPU NVLink 或 NVLink-C2C 的速度存取資源。
存取路徑：
- 本機（local）存取：走本地高頻寬 NVLink-C2C。
- 遠端（remote）存取：走 GPU NVLink，部分情況也用 NVLink-C2C。
透過 EGM，GPU thread 可跨 NVSwitch fabric 存取 CPU attached memory 與 HBM3 等全部記憶體資源。

Preliminaries：EGM Platforms 系統拓撲

目前 EGM 可在三種平台啟用：

平台	組成	互連
(1) Single-Node, Single-GPU	Arm-based CPU + CPU attached memory + 一顆 GPU	CPU 與 GPU 之間高頻寬 C2C（Chip-to-Chip）
(2) Single-Node, Multi-GPU	多顆 Arm-based CPU（各帶 attached memory）+ 多顆 GPU	GPU 經 NVLink-based network 連接
(3) Multi-Node, Multi-GPU	兩個以上的 (1) 或 (2) 單節點系統	節點間經 NVLink-based network 連接

(1) Single-Node Single-GPU      (2) Single-Node Multi-GPU
  +-----+   C2C   +-----+         CPU0-mem   CPU1-mem
  | CPU |=========| GPU |          |           |
  +-----+         +-----+         GPU0--NVLink--GPU1 ... (NVSwitch fabric)

(3) Multi-Node Multi-GPU
  [ Node A (1)/(2) ] ==NVLink network== [ Node B (1)/(2) ]

Warning

用 cgroups 限制可見裝置會阻斷 EGM 的路由並造成效能問題。請改用 CUDA_VISIBLE_DEVICES 來限制裝置。

Preliminaries：Socket Identifiers（NUMA node id）

NUMA（Non-Uniform Memory Access）把記憶體切成多個 node，每個 node 有自己的 processor 與 memory，並由 OS 指派唯一識別碼（numaID）。
EGM 使用 OS 指派的 NUMA node identifier 來表達記憶體放置位置。
重點：numaID 不等於 device 的 ordinal；它關聯到「最近的 host node」。
取得方式：用 cuDeviceGetAttribute 搭配 CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID。

int numaId;
cuDeviceGetAttribute(&numaId, CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID, deviceOrdinal);

傳入 device ordinal，回傳該裝置最近 host node 的 numaID，後續配置與授權都用它當 location id。

Preliminaries：Allocators 與 EGM 支援

把系統記憶體映射成 EGM 不會造成效能問題；存取遠端 socket 的系統記憶體（映射為 EGM）反而更快，因為流量保證走 NVLink。
目前支援兩種 allocator（搭配正確的 location type 與 NUMA 識別碼）：
- cuMemCreate：Virtual Memory Management（VMM）。
- cudaMemPoolCreate：Stream Ordered Memory Allocator（記憶體池）。

Preliminaries：對現有 API 的記憶體管理擴充

EGM 記憶體可用 VMM（cuMemCreate）或 Stream Ordered（cudaMemPoolCreate）映射。
使用者負責「配置實體記憶體」並把它「映射到所有 socket 的虛擬位址空間」。
為了讓這些 allocator 用 NUMA-like node id 理解配置位置，新增了 CUDA property type：

CUDA Type	用於
`CU_MEM_LOCATION_TYPE_HOST_NUMA`	`CUmemAllocationProp`（給 `cuMemCreate`）
`cudaMemLocationTypeHostNuma`	`cudaMemPoolProps`（給 `cudaMemPoolCreate`）

Important

Multi-node, multi-GPU 平台需要 interprocess communication（IPC）。EGM 同時建議搭配 Virtual Memory Management 與 Stream Ordered Memory Allocator 的章節一起理解。

Using EGM：Single-Node, Single-GPU

任何現有的 CUDA host allocator、以及 system allocated memory，都能直接享有高頻寬 C2C 的好處。
對使用者而言，「本機存取」就跟今天一個 host allocation 一樣，無需額外設定。
細部記憶體 allocator 與 page size 調校請參考 tuning guide。

Using EGM：Single-Node, Multi-GPU

多 GPU 系統中，使用者必須提供放置（placement）所需的 host 資訊；自然的表達方式就是 NUMA node id，EGM 正是採此做法。
流程：先用 cuDeviceGetAttribute 查出最近的 NUMA node id，再用 VMM API 或 CUDA Memory Pool 配置與管理 EGM 記憶體。

方法 A：使用 VMM API

建立實體記憶體後援（physical backing）。配置時須明確指定 CU_MEM_LOCATION_TYPE_HOST_NUMA 為 location type、numaID 為 location id，並對齊平台適當的 granularity：

CUmemAllocationProp prop{};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.location.type = CU_MEM_LOCATION_TYPE_HOST_NUMA;
prop.location.id = numaId;
size_t granularity = 0;
cuMemGetAllocationGranularity(&granularity, &prop, MEM_ALLOC_GRANULARITY_MINIMUM);
size_t padded_size = ROUND_UP(size, granularity);
CUmemGenericAllocationHandle allocHandle;
cuMemCreate(&allocHandle, padded_size, &prop, 0);

保留位址空間並映射到指標（此步驟無 EGM 專屬變更）：

CUdeviceptr dptr;
cuMemAddressReserve(&dptr, padded_size, 0, 0, 0);
cuMemMap(dptr, padded_size, 0, allocHandle, 0);

明確設定存取保護，否則存取映射空間會 crash。同樣以 HOST_NUMA + numaId，並為 host node 與 GPU 各建一個 access descriptor，賦予讀寫權限：

CUmemAccessDesc accessDesc[2]{{}};
accessDesc[0].location.type = CU_MEM_LOCATION_TYPE_HOST_NUMA;
accessDesc[0].location.id = numaId;
accessDesc[0].flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
accessDesc[1].location.type = CU_MEM_LOCATION_TYPE_DEVICE;
accessDesc[1].location.id = currentDev;
accessDesc[1].flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
cuMemSetAccess(dptr, size, accessDesc, 2);

重點：VMM 路徑的三步驟是 create（HOST_NUMA backing）→ reserve+map（無 EGM 變更）→ setAccess（必做，否則存取會 crash）。

方法 B：使用 CUDA Memory Pool

在 home device 上建立 memory pool，location type 用 cudaMemLocationTypeHostNuma、id 用 numaId：

cudaSetDevice(homeDevice);
cudaMemPoolProps props{};
props.allocType = cudaMemAllocationTypePinned;
props.location.type = cudaMemLocationTypeHostNuma;
props.location.id = numaId;
cudaMemPoolCreate(&memPool, &props);

用既有的 peer access API 給其他裝置直連存取（direct connect peer access）：

cudaMemAccessDesc desc{};
desc.flags = cudaMemAccessFlagsProtReadWrite;
desc.location.type = cudaMemLocationTypeDevice;
desc.location.id = accessingDevice;
cudaMemPoolSetAccess(memPool, &desc, 1);

將 pool 設給 resident device，再用 cudaMallocAsync 配置：

cudaDeviceSetMemPool(residentDevice, memPool);
cudaMallocAsync(&ptr, size, memPool, stream);

重點：Memory Pool 路徑用熟悉的 cudaMemPoolSetAccess 授權 peer，再以 stream-ordered 的 cudaMallocAsync 取得 EGM 指標。

Warning

EGM 以 2MB pages 映射。存取「非常大的配置」時可能遇到更多 TLB miss。

Using EGM：Multi-Node, Multi-GPU

除了記憶體配置外，remote peer access 沒有 EGM 專屬改動，遵循 CUDA inter-process（IPC）協定。
在來源節點（Node A）用 cuMemCreate 配置，除了 HOST_NUMA + numaID 之外，還須將 requestedHandleTypes 設為 CU_MEM_HANDLE_TYPE_FABRIC：

CUmemAllocationProp prop{};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_FABRIC;
prop.location.type = CU_MEM_LOCATION_TYPE_HOST_NUMA;
prop.location.id = numaId;
// ... cuMemGetAllocationGranularity / ROUND_UP / assert(padded_size % page_size == 0)
CUmemGenericAllocationHandle allocHandle;
cuMemCreate(&allocHandle, padded_size, &prop, 0);

Node A 匯出 fabric handle，透過 TCP/IP 傳給 Node B；Node B 匯入後即可如一般 fabric handle 使用：

// Node A: 匯出
cuMemExportToShareableHandle(&fabricHandle, allocHandle, CU_MEM_HANDLE_TYPE_FABRIC, 0);
// fabricHandle 經 TCP/IP 送到 Node B

// Node B: 匯入
cuMemImportFromShareableHandle(&allocHandle, &fabricHandle, CU_MEM_HANDLE_TYPE_FABRIC);

Node B 匯入後，照常 reserve 位址空間並 cuMemMap，最後對「每一顆本地 GPU」設定存取：

// 讓 Node B 上全部 8 顆本地 GPU 都能存取位於 Node A 的 EGM 記憶體
CUmemAccessDesc accessDesc[8];
for (int i = 0; i < 8; i++) {
  accessDesc[i].location.type = CU_MEM_LOCATION_TYPE_DEVICE;
  accessDesc[i].location.id = i;
  accessDesc[i].flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
}
cuMemSetAccess(dptr, size, accessDesc, 8);

重點：跨節點靠 CU_MEM_HANDLE_TYPE_FABRIC 與 export/import shareable handle 完成共享；映射後仍須對每顆本地 GPU 逐一 cuMemSetAccess 授權。

Node A                                   Node B
cuMemCreate(FABRIC, HOST_NUMA)           cuMemImportFromShareableHandle(FABRIC)
        |                                       |
cuMemExportToShareableHandle  --TCP/IP-->  reserve + cuMemMap
                                                |
                                         cuMemSetAccess(8 local GPUs, RW)

考試/測驗重點

問題	答案
EGM 靠哪個互連達成高頻寬本機存取	NVLink-C2C（Chip-to-Chip）
遠端存取走什麼路徑	GPU NVLink（部分情況也用 NVLink-C2C），可跨 NVSwitch fabric
EGM 用什麼識別記憶體放置位置	OS 指派的 NUMA node id（numaID），非 device ordinal
取得 numaID 的 API/屬性	cuDeviceGetAttribute + CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID
支援哪兩種 allocator	cuMemCreate（VMM）/ cudaMemPoolCreate（Stream Ordered）
VMM 的 location type 列舉值	CU_MEM_LOCATION_TYPE_HOST_NUMA
Runtime pool 的 location type 列舉值	cudaMemLocationTypeHostNuma
多節點額外要設的 handle type	CU_MEM_HANDLE_TYPE_FABRIC
跨節點如何共享 handle	cuMemExportToShareableHandle / cuMemImportFromShareableHandle（經 TCP/IP）
為何不可用 cgroups 限制裝置	會阻斷 EGM 路由並掉效能；應改用 CUDA_VISIBLE_DEVICES
EGM 的 page 大小	2MB pages，超大配置可能更多 TLB miss
VMM 映射後不 setAccess 會怎樣	存取映射空間會 crash
Single-Node Single-GPU 需要特別設定嗎	不需要；既有 host allocator / system memory 直接受惠於 C2C

重點總覽

概觀：EGM 是什麼

Preliminaries：EGM Platforms 系統拓撲

Preliminaries：Socket Identifiers（NUMA node id）

Preliminaries：Allocators 與 EGM 支援

Preliminaries：對現有 API 的記憶體管理擴充

Using EGM：Single-Node, Single-GPU

Using EGM：Single-Node, Multi-GPU

方法 A：使用 VMM API

方法 B：使用 CUDA Memory Pool

Using EGM：Multi-Node, Multi-GPU

考試/測驗重點

Related Notes