理解 VMM 最好的类比是 “送快递的地址 vs 实际房子”：

没有 VMM 的世界 (cudaMalloc):
  你租了一套房（分配物理显存） 
  + 同时给了你一个地址（虚拟地址）
  ↓
  这两件事绑死了：退租（free）→ 地址也没了
  想换一间大房子？原地址作废，所有朋友给你寄的东西都丢了

有了 VMM 的世界:
  第 1 步：租一套房（cuMemCreate）- 只有房子，没挂地址
  第 2 步：注册一个地址（cuMemAddressReserve）- 只有地址，地址下还没房子
  第 3 步：把房子关联到地址（cuMemMap）- 现在能用了
  第 4 步：发钥匙给你（cuMemSetAccess）- 允许你访问
  
  好处：可以暂时退租房子（cuMemUnmap），但地址保留
       以后你重新租房，挂回同一个地址
       你的朋友（Tensor、NCCL handle）一直用这个地址，不会丢

传统 cudaMalloc 的问题

1
2
3
4
5
6
7
8

// 传统方式 - 一步到位
float *ptr = (float*)cudaMalloc(10 GB);
// ptr 既是虚拟地址，也绑定了 10 GB 物理显存
// 这两者无法分离

cudaFree(ptr);
// 虚拟地址 和 物理显存 同时消失
// 所有存着 ptr 值的对象都变成悬挂指针 (dangling pointer)

具体例子展示问题所在：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12

float *weights = cudaMalloc(10 GB);    // 权重
nccl_comm->buffer = weights;            // NCCL 记住了这个地址

// 想临时释放 weights 给推理用？
cudaFree(weights);                      // 释放成功
// 但 nccl_comm->buffer 现在指向非法地址！
// 下次 NCCL 用这个 comm 就崩了

// 等推理完，想重新加载 weights
weights = cudaMalloc(10 GB);            // 新地址 (比如原来是 0x7f00, 新的是 0x9200)
// 现在 nccl_comm->buffer 还是旧地址 0x7f00！
// 必须把所有引用这个地址的对象全部更新，工作量巨大

VMM API 的四步分配

CUDA VMM API (CUDA 10.2+) 把这个过程拆成了独立的四步：

Step 1: 创建物理显存句柄（只分配物理）

1
2
3
4
5
6
7
8
9

CUmemGenericAllocationHandle handle;
CUmemAllocationProp prop = {0};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
prop.location.id = 0;  // GPU 0

cuMemCreate(&handle, size, &prop, 0);
// 此时：物理显存已分配，但还没法用 (没地址)
// handle 只是一个 "号码牌"，不是指针

Step 2: 预留虚拟地址空间（只占地址，不用物理）

1
2
3
4
5

CUdeviceptr dptr;
cuMemAddressReserve(&dptr, size, 0, 0, 0);
// 此时：虚拟地址 dptr 已预留，但地址下没有物理显存
// dptr = 0x7f0000000000 (假设)
// 访问 dptr 会 segfault

Step 3: 把物理显存映射到虚拟地址

1
2
3
4
5

cuMemMap(dptr, size, 0, handle, 0);

// 此时：虚拟地址 dptr 和物理显存 handle 建立了映射关系

// 但还没权限访问

Step 4: 设置访问权限

1
2
3
4
5
6
7

CUmemAccessDesc access_desc = {0};
access_desc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
access_desc.location.id = 0;
access_desc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;

cuMemSetAccess(dptr, size, &access_desc, 1);
// 现在 dptr 可以像 cudaMalloc 返回的指针一样使用

关键的"解耦"操作

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16

// 暂时释放物理显存，但保留虚拟地址
cuMemUnmap(dptr, size);        // 物理显存回到 GPU 空闲池
cuMemRelease(handle);          // 释放物理显存句柄

// 此时状态：
// - dptr 虚拟地址仍然存在（cuMemAddressReserve 时预留的）
// - 访问 dptr 会 segfault（没映射到物理显存）
// - 但 dptr 这个值没变！所有记住它的对象都还知道该找谁

// 想用了，重新分配物理显存，映射到同一个 dptr
CUmemGenericAllocationHandle new_handle;
cuMemCreate(&new_handle, size, &prop, 0);
cuMemMap(dptr, size, 0, new_handle, 0);  // ← 同一个 dptr！
cuMemSetAccess(dptr, size, &access_desc, 1);
// 现在 dptr 又可以用了，指向新的物理显存
// 所有记住 dptr 的对象不需要更新！