GPU 如何访问另一块 GPU 的显存：同进程 P2P、跨进程 IPC 与 VMM Fabric

CUDA 侧按 CUDA 13.4 Runtime / Driver API 与 Programming Guide 核对；SGLang 侧按 upstream main928e285b（2026-11-01）核对。P2P、cudaIpc、VMM POSIX FD 有同机实测数据；PCIe P2P、ACS/IOMMU 配置变化、FABRIC/IMEX 与远端 attention 未测。文中带宽数字区分规范值、仓库测试和补充测量。

八张 GPU 各跑一个推理副本，同一段长 system prompt 会在每张卡上生成一份相同的 KV。八份副本占八份显存，一个自然的想法是让 GPU 0 直接读 GPU 1 的 KV，而不是再复制一份。

但要让 GPU 0 直接读，先得解决地址问题。GPU 1 的 cudaMalloc 返回的是它所在 context 里的地址，不能直接当作 GPU 0 或另一个进程可用的指针。同一进程里，UVA 会让地址看起来在同一个空间，GPU 0 要真正解引用仍需要 peer access；跨进程时，地址空间各自独立，必须先传句柄、重新建立映射。

cudaMemcpyPeer 能完成拷贝，但它没有回答 GPU 0 如何拿到一个可解引用的 GPU 1 地址。本文先按同一进程、跨进程、跨节点三条边界拆地址机制，再看 kernel 访问、同步和拓扑；最后用两个实现回答另一个问题：地址可达之后，哪些数据真的适合直接读。

一、先分清：拷贝、映射和直接解引用

“GPU 访问另一块 GPU 的显存”常被说成一个动作，实际至少有三层。

拷贝不改变指针可见性，映射不搬数据，解引用要求映射和访问权限都已就绪。三种动作的失败方式不同：拷贝是 API 错误，映射是句柄或权限问题，解引用是地址在当前 context 里无效。

图 1：同进程靠 UVA 和 peer access；跨进程靠句柄导入；跨节点要靠 fabric 或网卡中转。

二、同一进程：UVA 与 peer access

2.1 一个指针为什么在另一张卡上无效

CUDA 给每张 GPU 维护独立的 primary context，里面有独立的分配堆、stream 池和 context 状态。cudaSetDevice(1) 之后拿到的 float *d1，只在 GPU 1 的 context 里有意义。切回 GPU 0 后，如果没有启用 peer access，这个指针不能直接解引用。

同一进程内，64 位 Linux 上多张 GPU 通常共享统一虚拟地址空间（UVA）；地址存在不等于可访问，GPU 0 读 GPU 1 的 HBM 还需要 peer access。

2.2 查询和启用 peer access

cudaDeviceCanAccessPeer() 查询一个方向是否支持直接访问：

   
   
   
   int can = 0;
cudaDeviceCanAccessPeer(&can, 0, 1);  // GPU 0 能否直接访问 GPU 1

返回 1 表示硬件和驱动允许，访问授权仍要调用 cudaDeviceEnablePeerAccess()。这个调用是单向的；要让 GPU 1 也能访问 GPU 0，需要反过来再调一次。授权成功后，peer 设备的所有 allocation 立刻可被当前设备访问，直到显式 disable 或 device reset。重复启用会返回 cudaErrorPeerAccessAlreadyEnabled。

   
   
   
   cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1, 0);

// d1 是在 GPU 1 上通过 cudaMalloc 分配的指针
kernel<<<blocks, threads>>>(d1);

这类授权是 per-device 的，不是 per-allocation；通常在初始化时对所有 GPU 对一次做完，不要每次拷贝前重复调用。

2.3 kernel 直接读，和 copy 路径

启用 peer access 后，GPU 0 的 kernel 可以直接解引用 d1。CUDA 运行时和驱动负责让这段地址在当前设备上可访问，load/store 请求通过 NVLink 或 PCIe 发到目标 GPU。CUDA 文档把这类访问描述为 fine-grained peer-to-peer GPU load/store。

另一条路径是拷贝：

   
   
   
   cudaMemcpyPeerAsync(dst_on_gpu1, 1, src_on_gpu0, 0, bytes, stream);

cudaMemcpyPeer 异步于 host，但会与 current device、srcDevice、dstDevice 上已排队和后续提交的异步工作串行化。cudaMemcpyPeerAsync 可以避开这种同步。CUDA 文档说 P2P 拷贝可以利用专门的 copy engine 和 NVLink 硬件，但没有保证每条路径一定走哪个引擎。

小块、稀疏、需要即时读取的数据适合 kernel 直接 load/store；大块、连续、可以预取的数据适合拷贝。两者都要求拓扑支持，区别是访问暴露给 kernel，还是先把数据搬到本地显存。

图 2：peer access 建立的是地址翻译权限；cudaMemcpyPeerAsync 走的是拷贝路径。

三、跨进程：cudaIpc 的句柄与生命周期

3.1 指针不能直接传

两个进程各有独立的虚拟地址空间。进程 A 的 GPU 指针 0x7f... 在进程 B 里没有意义，哪怕它们在同一张 GPU 上。跨进程共享要传一个句柄，由运行时在导入进程里重新建立映射。

CUDA Runtime 提供了 cudaIpcGetMemHandle / cudaIpcOpenMemHandle：

   
   
   
   // 导出进程
cudaIpcMemHandle_t handle;
cudaIpcGetMemHandle(&handle, d_ptr);
// 通过 Unix socket、共享内存或其他 IPC 把 handle 交给对方

// 导入进程
void *remote = NULL;
cudaIpcOpenMemHandle(&remote, handle, cudaIpcMemLazyEnablePeerAccess);
kernel<<<blocks, threads>>>(remote);
cudaDeviceSynchronize();  // 等 kernel 用完 remote
cudaIpcCloseMemHandle(remote);

cudaIpcOpenMemHandle 把另一个进程导出的内存映射到当前进程的 device 地址空间。对不同 device 的 context，它会尝试像 cudaDeviceEnablePeerAccess 一样建立 peer access；是否可行仍要看 cudaDeviceCanAccessPeer。

3.2 句柄不是指针，也不是永久有效

cudaIpcGetMemHandle 接受由 cudaMalloc 分配的 device memory 基地址。它不是通用的“导出任意 GPU 内存”接口：

• • cudaMallocManaged 分配不支持 IPC。官方 Programming Guide 写得很直接：IPC API 不支持 managed allocation。
• • 句柄面向同一台主机上的进程，不是跨节点机制。跨节点要么用 VMM 的 FABRIC handle，要么走 RDMA/NCCL。
• • 同一个 handle 在当前 context 重复打开会增加引用计数，返回已有指针；导入方必须用 cudaIpcCloseMemHandle 释放导入映射。
• • 导出进程在导入方 close 之前对导出区域调用 cudaFree，是未定义行为。
• • 同一个进程里，每张 device 的 IPC handle 在另一个进程中只能由一个 context 打开；多 context 共享要自己协调。
• • 导入后的地址不保证和导出进程相同，不能把地址值当成跨进程标识。

句柄本身是一个能力凭证。拿到它的人可以在自己的进程里映射这块显存，所以传递句柄的通道要有访问控制。cudaIpc 解决的是地址映射，不解决权限模型。

3.3 cudaIpc 适合什么

它适合同一台机器上的多进程推理、前后端分离、共享一块固定的权重或缓存。它的 API 少，路径短，但分配来源被限制在 cudaMalloc，生命周期也要由应用自己协调。

如果分配来自 CUDA VMM 或新的 allocator，cudaIpc 不一定能覆盖。这时要看下一节的 shareable handle。

图 3：导出的是句柄；导入方负责映射、授权和释放。

四、VMM：POSIX FD 与 Fabric

4.1 先把 allocation 和 mapping 分开

CUDA VMM 把一次分配拆成几步：

1. 1. cuMemAddressReserve 保留一段虚拟地址。
2. 2. cuMemCreate 创建物理 allocation，此时没有任何 device 或 host 映射。
3. 3. cuMemExportToShareableHandle 把 allocation 导出成可传递的句柄。
4. 4. 接收方 cuMemImportFromShareableHandle 导入句柄。
5. 5. cuMemMap 把 allocation 映射到已保留的虚拟地址。
6. 6. cuMemSetAccess 授予目标设备访问权限。

简化示意：

   
   
   
   CUmemAllocationProp prop = {0};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
prop.location.id = dev;
prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_FABRIC;

CUmemGenericAllocationHandle alloc;
cuMemCreate(&alloc, size, &prop, 0);

CUmemFabricHandle shareable;
cuMemExportToShareableHandle(&shareable, alloc, CU_MEM_HANDLE_TYPE_FABRIC, 0);

// 接收进程
CUmemGenericAllocationHandle imported;
cuMemImportFromShareableHandle(&imported, &shareable, CU_MEM_HANDLE_TYPE_FABRIC);

CUdeviceptr va;
cuMemAddressReserve(&va, size, 0, 0, 0);
cuMemMap(va, size, 0, imported, 0);

CUmemAccessDesc access = {0};
access.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
access.location.id = dev;
access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
cuMemSetAccess(va, size, &access, 1);

CU_MEM_HANDLE_TYPE_FABRIC 的输出是 64 字节的 CUmemFabricHandle 结构体，不是指针；POSIX FD 分支则用 int 文件描述符。

这段代码省略了错误处理、释放顺序和 granularity 计算。实际使用时，cuMemCreate 的 size 必须是 cuMemGetAllocationGranularity 按 CU_MEM_ALLOC_GRANULARITY_MINIMUM 返回值的整数倍，cuMemMap 的地址、大小和偏移也要按同一粒度对齐。

cuMemMap 只建立映射，cuMemSetAccess 才授权；导入方必须自己完成这两步，不能假设 handle 导入后就得到可读写指针。示例省略了 unmap/release；跨设备导入时，access.location.id 要用导入方设备，不能沿用导出方的 dev。

4.2 POSIX FD 和 FABRIC 的边界

VMM 的 handle 类型决定共享范围：

FABRIC handle 不是“有 NVLink 就能用”。CUDA 文档要求导出和导入进程能访问同一个 IMEX channel；没有同一 channel 时，导入会返回 CUDA_ERROR_NOT_PERMITTED。IMEX 的权限按用户组织，多租户环境需要单独划分 channel。

VMM 的 host allocation 不支持 IPC。CU_MEM_LOCATION_TYPE_HOST 的 allocation 必须把 requestedHandleTypes 设为 0，设成其他值会返回 CUDA_ERROR_INVALID_VALUE。这条边界常被忽略：能通过 VMM 分配的 host 内存，不等于能通过 VMM 跨进程共享。

4.3 cudaIpc 与 VMM 怎么选

还有第三条路径：cudaMallocAsync 的 memory pool。IPC-capable pool 可以用 cudaMemPoolExportToShareableHandle / cudaMemPoolImportFromShareableHandle 共享 pool，再用 cudaMemPoolExportPointer / cudaMemPoolImportPointer 共享具体指针。它共享的是 allocator 的 pool，不是逐块 cudaMalloc；创建 pool 时 handle type 不能是 None。

新代码如果要在 NVLink fabric 内跨节点，或者需要把 allocation 和 mapping 分开管理，VMM 更合适；没有 fabric 的跨节点仍走 RDMA/NCCL。cudaIpc 更简单，边界也更窄。

五、kernel 视角：FlashAttention 怎么读 KV