先分清“可访问”“驻留”和“同步”

CUDA 内存问题常因把三个概念混在一起而出错：

• 可访问性：某个 processor/context 是否拥有读写该地址的权限；
• 驻留位置：对应物理页当前位于 GPU HBM、host DRAM 还是其他位置；
• 一致性与顺序：一个执行者的写入何时对另一个执行者可见。

统一虚拟地址让 pointer 值更统一，Unified Memory 让页面可迁移，system memory 能否被 GPU 直接访问则取决于平台。它们都不意味着“CPU 和 GPU 可以无同步地同时读写，并且永远拥有本地带宽”。

内存管理方案总览

这些方案可以组合。例如：应用用 VMM 建大型地址空间，框架在其上构建 allocator；或业务数据用 managed memory，热点工作区由 stream-ordered pool 分配。

Unified Memory 的页迁移模型

cudaMallocManaged 返回 CPU/GPU 均可寻址的 managed pointer。典型离散 GPU 系统上，首次触碰或 ownership 变化可能触发 page fault 和迁移：

时刻 t0：CPU 初始化 pages   [Host DRAM]
                          │ GPU 首次访问，page fault
                          ▼
时刻 t1：页面迁移/映射     [GPU Memory]
                          │ CPU 再次随机访问
                          ▼
时刻 t2：迁回或远程访问     [Host / GPU，取决于平台与策略]

页错误批次、页面大小、互连和访问模式共同决定成本。第一次 kernel 慢而后续变快，经常只是迁移成本被摊到第一次访问，不应误判为 JIT 或 cache 的全部影响。

Managed Memory 最小用法

float* p = nullptr;
std::size_t bytes = n * sizeof(float);
CUDA_CHECK(cudaMallocManaged(&p, bytes));

for (std::size_t i = 0; i < n; ++i) p[i] = static_cast<float>(i);

kernel<<<grid, block>>>(p, n);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());  // CPU 读取结果前建立顺序

consume_on_cpu(p, n);
CUDA_CHECK(cudaFree(p));

基础正确性仍要求同步。某些平台支持更强的 concurrent managed access 或硬件一致性，但可用能力必须查询，不能把一种系统的经验推广到所有系统。

Prefetch 与内存建议

若访问阶段清晰，可提前表达意图：

int dev = 0;
CUDA_CHECK(cudaMemPrefetchAsync(p, bytes, dev, stream));
kernel<<<grid, block, 0, stream>>>(p, n);

// CPU 即将集中读取时，再预取回 CPU。
CUDA_CHECK(cudaMemPrefetchAsync(p, bytes, cudaCpuDeviceId, stream));
CUDA_CHECK(cudaStreamSynchronize(stream));

内存 advice 可表达 preferred location、read mostly、accessed by 等偏好。它们是策略提示，不是强制放置或同步操作。

访问模式稳定且分阶段？ → prefetch 通常有效
多 GPU 只读共享？       → read-mostly + accessed-by 值得评估
CPU/GPU 反复写同一页？  → 先改变数据分区，避免 ping-pong
数据大于显存？           → 分块/流式算法优先，UM 超额订阅需实测

Unified Memory 性能反模式

• CPU 与 GPU 交替写相邻字段，导致同一页面反复迁移；
• 多 GPU 随机访问同一批可写 managed pages；
• 首次访问成本未计入端到端 benchmark，却拿稳态 kernel 时间做结论；
• 将 managed pointer 当作自动预取器，完全不考虑访问阶段；
• 超额订阅时工作集没有局部性，造成持续 eviction/fault；
• 在不同平台上假定相同页面、一致性或原子能力。

优化的首要动作通常是重新划分 ownership，让一个阶段的大部分页面由同一个 processor 连续使用。

System-allocated Memory 与平台差异

现代平台可能允许 GPU 访问 malloc/系统分配的内存，具体能力与 HMM、ATS、页表架构、CPU-GPU 互连及操作系统有关。Grace Hopper 等硬件一致平台与 PCIe 离散 GPU 的性能特征不同。

运行时查询比型号推断可靠：

• managed memory 与 concurrent managed access；
• pageable memory access / host page table 能力；
• direct managed access from host；
• memory pools、VMM 和 handle type 支持；
• P2P、原子与 fabric 能力。

即使能访问 system memory，大带宽重复访问仍可能值得显式迁入 GPU 本地内存。

为什么 cudaMallocAsync 很重要

传统 cudaMalloc/cudaFree 可能引入设备范围的昂贵协调。Stream-ordered allocator 把“何时可用/何时可复用”纳入 stream dependency：

void* tmp = nullptr;
CUDA_CHECK(cudaMallocAsync(&tmp, bytes, stream));
stage1<<<g1, b1, 0, stream>>>(tmp);
stage2<<<g2, b2, 0, stream>>>(tmp);
CUDA_CHECK(cudaFreeAsync(tmp, stream));

stream: [alloc becomes usable][stage1][stage2][free becomes reusable]
pool:   └──────────── backing allocation retained/recycled ─────────┘

API 返回 pointer 不代表其他 stream 可以立刻访问。跨 stream 使用需要 event/order 关系；free 后也不能再由无序的工作访问。

Memory Pool

每个 device 有默认 pool，也可创建显式 pool：

cudaMemPoolProps props{};
props.allocType = cudaMemAllocationTypePinned;
props.location.type = cudaMemLocationTypeDevice;
props.location.id = device;

cudaMemPool_t pool{};
CUDA_CHECK(cudaMemPoolCreate(&pool, &props));

std::uint64_t threshold = 512ull << 20;
CUDA_CHECK(cudaMemPoolSetAttribute(
    pool, cudaMemPoolAttrReleaseThreshold, &threshold));

void* p = nullptr;
CUDA_CHECK(cudaMallocFromPoolAsync(&p, bytes, pool, stream));

重要属性包括 reserved/used current 与 high watermark、release threshold，以及跨 stream 的复用策略。池保留内存会让 nvidia-smi 看见较高占用，这不等于应用泄漏；要同时检查 pool 的 used 与 reserved。

跨 Stream 复用证明

stream A: [use p]──[freeAsync p]──record(E)
                                  │
stream B: ─────────────────────wait(E)──[alloc/reuse]

pool 可根据 event dependency、机会式规则或内部依赖决定复用。内部自动依赖可能引入意外串行化；对关键路径，显式 event 往往更容易解释和分析。

Graph capture 中的异步分配会形成 allocation/free node，其 ownership 和普通 pool allocation 不完全相同。Graph 生命周期、更新和重复实例化要单独测试。

Runtime allocator 与 VMM 的抽象差异

cudaMalloc / cudaMallocAsync
    一次调用同时取得“地址 + 物理 backing + 访问权”

Driver VMM
    reserve virtual address
            + create physical allocation
            + map VA → physical
            + set per-device access

VMM 将虚拟地址与物理存储拆开，类似操作系统的 mmap 思想。这使稳定地址、按需扩容、稀疏映射、精细 P2P 权限和跨进程共享成为可能，但 API 更底层。

VMM 基本流水线

1 query granularity/support
2 cuMemAddressReserve         → 仅保留 VA
3 cuMemCreate                 → 创建物理 allocation handle
4 cuMemMap                    → 将物理内存映射到 VA
5 cuMemSetAccess              → 授予某 device 读写权限
6 kernel/memcpy 使用普通 pointer
7 cuMemUnmap
8 cuMemRelease                → 释放物理 handle
9 cuMemAddressFree            → 释放 VA

精简代码骨架：

CUmemAllocationProp prop{};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
prop.location.id = device;

size_t granularity = 0;
CU_CHECK(cuMemGetAllocationGranularity(
    &granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM));
size_t mapped_bytes = round_up(bytes, granularity);

CUdeviceptr va = 0;
CU_CHECK(cuMemAddressReserve(&va, mapped_bytes, 0, 0, 0));

CUmemGenericAllocationHandle handle{};
CU_CHECK(cuMemCreate(&handle, mapped_bytes, &prop, 0));
CU_CHECK(cuMemMap(va, mapped_bytes, 0, handle, 0));

CUmemAccessDesc access{};
access.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
access.location.id = device;
access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
CU_CHECK(cuMemSetAccess(va, mapped_bytes, &access, 1));

cuMemMap 后若未 cuMemSetAccess，kernel 访问仍会失败。所有 size/offset 必须满足查询到的 granularity，而不是写死 64 KiB 或 2 MiB。

VMM 能解决哪些高级问题

稳定虚拟地址下扩容

先保留大 VA 区间，只为当前容量创建并映射物理页；增长时映射后续区间。这样上层保存的 base pointer 不变，适合大型 KV cache、数据库或 arena。

reserved VA: [mapped A][mapped B][unmapped..............]
grow:        [mapped A][mapped B][new mapping C][.......]

稀疏和按需 backing

逻辑地址空间可以大于当前物理 allocation，仅映射实际需要的区间。未映射区域绝不能被 kernel 触碰；应用自己负责 residency 元数据和错误边界。

精细多 GPU 权限

传统 cudaDeviceEnablePeerAccess 会让一类 allocation 广泛映射到 peer。VMM 可以只把需要共享的 allocation 映射到指定 GPU，并分别设置权限，减少地址映射开销和权限面。

IPC 与 Fabric handle

物理 allocation handle 可以导出为操作系统句柄，在另一个进程导入后映射。支持的平台还可使用 fabric handle 跨节点共享。句柄只描述 GPU allocation，实际传输 handle、认证、生命周期与错误恢复仍需 IPC/MPI 等通道。

IPC 生命周期图

Process A                                  Process B
cuMemCreate
    │
export shareable handle ── socket/MPI ──→ import handle
    │                                          │
reserve VA + map + access                 reserve VA + map + access
    │                                          │
GPU work A  ←──── event/protocol/order ───→ GPU work B
    │                                          │
unmap/release only after all users have stopped and acknowledged

“能映射”不提供进程间同步。共享数据结构仍需要协议、内存顺序和崩溃恢复设计。OS handle 和 CUDA handle 的所有权也必须分别关闭。

VMM 不适合什么

• 只需要几个固定 buffer 的普通应用；
• 团队无法维护严格的映射/权限/释放状态机；
• 可由 NCCL、NVSHMEM 或框架 allocator 直接解决的问题；
• 误以为 VMM 会自动迁移、自动负载均衡或自动同步；
• 未查询 feature/granularity 就按某款 GPU 写死实现。

内存问题诊断表

选型流程

固定、长寿命 GPU buffer？ ─────────────→ cudaMalloc / RAII
频繁创建临时工作区？ ─────────────────→ cudaMallocAsync + pool
CPU/GPU 共享复杂结构，迁移可接受？ ───→ Unified Memory + prefetch/advice
I/O staging 或异步 H2D/D2H？ ─────────→ pinned host memory
稳定 VA、稀疏映射、精细 P2P/IPC？ ───→ Driver VMM
标准 collective/PGAS 已覆盖？ ────────→ NCCL/NVSHMEM，避免自建 VMM 协议

练习

1. 对同一数组分别测试 cudaMalloc、managed on-demand、managed prefetch，分开记录首轮与稳态。
2. 用 cudaMallocAsync 构造两个 stream 的生产消费，分别用 event 和 device sync 建依赖，比较时间线。
3. 查询默认 pool 的 used/reserved high watermark，测试不同 release threshold。
4. 用 VMM 保留 1 GiB VA，但只映射前 64 MiB；扩容一次并保证 base pointer 不变。

官方参考

• CUDA Programming Guide: Unified Memory
• CUDA Programming Guide: Stream-Ordered Memory Allocator
• CUDA Programming Guide: Interprocess Communication
• CUDA Programming Guide: Virtual Memory Management
• CUDA Runtime API: Memory Pools
• CUDA Driver API: Virtual Memory Management