大模型架构

Appearance

CUDA 跨进程共享显存 ​

在单机多 GPU 程序中，不同进程可以把一块 GPU Device Memory 映射到各自的 CUDA 虚拟地址空间，再由 Kernel 通过普通指针直接访问。NVIDIA 提供两类进程间显存共享接口：Legacy CUDA IPC 和 VMM-based IPC。前者把 cudaMalloc Allocation 作为整体导出，后者使用 CUDA Virtual Memory Management（VMM）接口，把物理 Allocation、虚拟地址和访问权限拆开管理。

相关术语定义如下：

  • Legacy CUDA IPC：cudaIpcGetMemHandle、cudaIpcOpenMemHandle 及对应的 cuIpc* Memory API；同一 API 家族还提供 IPC Event。
  • VMM-based IPC：特指 cuMemExportToShareableHandle、cuMemImportFromShareableHandle 和 cuMemMap 等 VMM Driver API 组成的共享路径。
  • CUDA 进程间通信：Legacy CUDA IPC、VMM-based IPC 等进程间 CUDA 资源共享机制的总称。

跨进程共享显存不等同于 CUDA C++ 的 __shared__ Memory：

名称所在位置可见范围主要用途
CUDA Shared MemorySM 片上存储同一 Thread BlockBlock 内线程协作
跨进程共享 Device MemoryHBM/GDDR，经 L2 和 GPU Interconnect 访问建立映射的 CUDA ContextGPU 间或进程间交换数据

以下内容限定为 Linux 单机进程间共享。跨节点通信还需要 GPUDirect RDMA、NVSHMEM、NCCL 或 CUDA Fabric Handle 等机制。

1. 地址映射与数据传输 ​

CUDA Device Pointer 是当前进程 CUDA 虚拟地址空间中的地址，不是可以直接发送给其他进程使用的物理地址：

text
Process 0 / CUDA Context 0
  pointer_a ──VA Mapping──┐
                          ├── GPU 0 Physical Allocation
Process 1 / CUDA Context 1│
  pointer_b ──VA Mapping──┘

pointer_a 与 pointer_b 的数值可以不同，但它们映射到同一个底层 Allocation。Process 1 的 Kernel 对 pointer_b 执行 Load/Store 时，数据根据设备拓扑经过 NVLink、NVSwitch 或 PCIe 到达 GPU 0 显存。

因此，跨进程共享显存需要区分三层：

text
Allocation 与地址映射
  Legacy CUDA IPC 或 VMM-based IPC
          ↓ 提供可解引用的 Peer Pointer
数据访问
  Kernel Load/Store、Memcpy 或通信 Kernel
          ↓ 经过物理链路
同步与可见性
  CUDA Event、Stream Dependency、Barrier 或 CPU 控制通道

Legacy CUDA IPC 与 VMM-based IPC 只建立映射，不会自动复制数据，也不定义 All-Reduce、All-Gather 等通信语义。应用仍要决定谁读、谁写、何时可以复用缓冲区。

2. 使用条件 ​

Legacy CUDA IPC 和单机 VMM-based IPC 通常需要满足以下条件：

  1. 进程运行在同一台 Linux 主机上，并由兼容的 NVIDIA Driver 管理。
  2. GPU 支持 Unified Virtual Addressing；VMM 还要求 CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED。
  3. 跨 GPU 直接访问前，使用 cudaDeviceCanAccessPeer 或 cuDeviceCanAccessPeer 检查拓扑是否支持 P2P。
  4. Exporter 在所有 Importer 释放映射前保持原始 Allocation 存活。
  5. 程序在 fork() 前不要初始化 CUDA。更常见的生产方式是直接启动独立 Worker 进程，再通过 Unix Domain Socket 或进程组交换 Handle。

即使 P2P 可用，链路也不一定是 NVLink。PCIe P2P 与 NVLink 的带宽、延迟和原子操作能力不同，应通过 nvidia-smi topo -m 和实际 Profile 确认物理路径。

3. Legacy CUDA IPC ​

Legacy CUDA IPC 通常使用 CUDA Runtime API，把 cudaMalloc 创建的完整 Device Allocation 导出为 cudaIpcMemHandle_t：

text
Exporter                                      Importer

cudaMalloc                                    接收 cudaIpcMemHandle_t
    ↓                                                  ↓
cudaIpcGetMemHandle                          cudaIpcOpenMemHandle
    ↓                                                  ↓
通过 Unix Socket/Shared Memory 发送 Handle     获得本进程 Peer Pointer
                                                       ↓
                                              Kernel 直接访问远端显存
                                                       ↓
                                              cudaIpcCloseMemHandle
    ↓
等待 Importer 完成后 cudaFree

3.1 Exporter ​

Exporter 在 GPU 0 分配显存、完成初始化并导出 Handle。send_bytes 代表任意可靠的进程间控制通道；它传递的只是 Handle 和元数据，不传递 Device Memory 内容。

cpp
#include <cuda_runtime.h>

#include <cstddef>

struct IpcAllocationDesc {
    cudaIpcMemHandle_t handle;
    size_t bytes;
};

IpcAllocationDesc export_allocation(void** pointer, size_t bytes) {
    cudaSetDevice(0);
    cudaMalloc(pointer, bytes);

    // 在导出前完成需要对 Importer 可见的初始化。
    cudaMemset(*pointer, 0, bytes);
    cudaDeviceSynchronize();

    IpcAllocationDesc desc{};
    desc.bytes = bytes;
    cudaIpcGetMemHandle(&desc.handle, *pointer);
    return desc;
}

void exporter(int control_socket) {
    void* owner_pointer = nullptr;
    IpcAllocationDesc desc =
        export_allocation(&owner_pointer, 4 * 1024 * 1024);

    send_bytes(control_socket, &desc, sizeof(desc));

    // 等待 Importer 完成 Kernel 并关闭映射。
    wait_done(control_socket);

    cudaFree(owner_pointer);
}

cudaIpcGetMemHandle 要求传入 Allocation Base Pointer。若 Tensor 只是较大 Allocation 中的一个切片，应额外发送相对 Base 的 Offset，由 Importer 在打开 Base 后计算真实地址：

cpp
void* tensor_pointer = static_cast<char*>(base_pointer) + tensor_offset;

3.2 Importer ​

Importer 在 GPU 1 上打开 Handle。返回的 peer_pointer 属于 Importer 的 CUDA 地址空间，但底层仍是 Exporter 在 GPU 0 上创建的物理 Allocation。

cpp
__global__ void add_one(float* data, size_t count) {
    size_t index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < count) {
        data[index] += 1.0f;
    }
}

void importer(int control_socket) {
    IpcAllocationDesc desc{};
    recv_bytes(control_socket, &desc, sizeof(desc));

    cudaSetDevice(1);

    int can_access = 0;
    cudaDeviceCanAccessPeer(&can_access, 1, 0);
    if (!can_access) {
        throw std::runtime_error("GPU 1 cannot access GPU 0 through P2P");
    }

    void* peer_pointer = nullptr;
    cudaIpcOpenMemHandle(
        &peer_pointer,
        desc.handle,
        cudaIpcMemLazyEnablePeerAccess);

    size_t count = desc.bytes / sizeof(float);
    add_one<<<(count + 255) / 256, 256>>>(
        static_cast<float*>(peer_pointer), count);
    cudaDeviceSynchronize();

    // Open 得到的 Pointer 必须 Close，不能 cudaFree。
    cudaIpcCloseMemHandle(peer_pointer);
    notify_done(control_socket);
}

cudaIpcMemLazyEnablePeerAccess 允许 CUDA 在打开 Handle 时按需建立设备间 Peer Access。Importer 不应假设打开后的指针与 Exporter 原始指针数值相同。

3.3 Legacy CUDA IPC Event ​

内存映射本身不表示数据已经生产完成。一次性生产者—消费者流程可以额外共享 CUDA Event：

cpp
// Exporter
cudaEvent_t ready;
cudaEventCreateWithFlags(
    &ready,
    cudaEventDisableTiming | cudaEventInterprocess);

produce<<<grid, block, 0, producer_stream>>>(owner_pointer);
cudaEventRecord(ready, producer_stream);

cudaIpcEventHandle_t event_handle;
cudaIpcGetEventHandle(&event_handle, ready);
send_bytes(control_socket, &event_handle, sizeof(event_handle));

Importer 打开 Event，并让消费 Stream 等待生产 Stream：

cpp
// Importer
cudaIpcEventHandle_t event_handle;
recv_bytes(control_socket, &event_handle, sizeof(event_handle));

cudaEvent_t ready;
cudaIpcOpenEventHandle(&ready, event_handle);
cudaStreamWaitEvent(consumer_stream, ready, 0);

consume<<<grid, block, 0, consumer_stream>>>(peer_pointer);

共享 Event 适合表达 Stream 间依赖，但循环通信仍需设计 Buffer Slot、Epoch 或双向完成通知，避免生产者覆盖消费者尚未读完的数据。

4. VMM-based IPC ​

VMM 使用 CUDA Driver API，把传统 cudaMalloc 隐含完成的操作拆成四步：

text
cuMemCreate           创建物理 Allocation，尚无可访问 Pointer
cuMemAddressReserve   保留一段 GPU 虚拟地址
cuMemMap              建立 VA → Physical Allocation 映射
cuMemSetAccess        设置哪些 GPU 可以读写这段映射

同一物理 Allocation 可以被导出到其他进程，再分别映射到各自的虚拟地址空间：

text
Exporter                                      Importer

cuMemCreate                                   接收 POSIX fd
    ↓                                                   ↓
cuMemAddressReserve                           cuMemImportFromShareableHandle
    ↓                                                   ↓
cuMemMap                                      cuMemAddressReserve
    ↓                                                   ↓
cuMemSetAccess                                cuMemMap
    ↓                                                   ↓
cuMemExportToShareableHandle                  cuMemSetAccess
    ↓                                                   ↓
SCM_RIGHTS 发送 POSIX fd                      获得可访问的 Peer VA

4.1 创建可导出的 VMM Allocation ​

cuMemCreate 的 Size 必须按设备要求的 Granularity 向上对齐。VMM API 中的 CU_MEM_ALLOCATION_TYPE_PINNED 表示当前 Device Allocation 类型，不是 cudaHostAlloc 创建的 Host Pinned Memory。

cpp
#include <cuda.h>

#include <cstddef>

struct VmmAllocation {
    CUmemGenericAllocationHandle handle;
    CUdeviceptr address;
    size_t allocation_size;
};

size_t round_up(size_t value, size_t alignment) {
    return (value + alignment - 1) / alignment * alignment;
}

void set_rw_access(CUdeviceptr address, size_t size, int device_id) {
    CUmemAccessDesc access{};
    access.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
    access.location.id = device_id;
    access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
    cuMemSetAccess(address, size, &access, 1);
}

VmmAllocation create_exportable_allocation(
    int device_id,
    size_t requested_size) {
    CUmemAllocationProp prop{};
    prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
    prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
    prop.location.id = device_id;
    prop.requestedHandleTypes =
        CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;

    size_t granularity = 0;
    cuMemGetAllocationGranularity(
        &granularity,
        &prop,
        CU_MEM_ALLOC_GRANULARITY_MINIMUM);

    VmmAllocation allocation{};
    allocation.allocation_size =
        round_up(requested_size, granularity);

    // 此时只获得物理 Allocation Handle，尚不能读写。
    cuMemCreate(
        &allocation.handle,
        allocation.allocation_size,
        &prop,
        0);

    cuMemAddressReserve(
        &allocation.address,
        allocation.allocation_size,
        granularity,
        0,
        0);

    cuMemMap(
        allocation.address,
        allocation.allocation_size,
        0,
        allocation.handle,
        0);

    // cuMemMap 不会自动授予访问权限。
    set_rw_access(
        allocation.address,
        allocation.allocation_size,
        device_id);

    return allocation;
}

创建 Allocation 前应检查 VMM 能力：

cpp
int supported = 0;
cuDeviceGetAttribute(
    &supported,
    CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED,
    device);

4.2 导出 POSIX fd ​

Exporter 将 VMM Allocation 导出为 POSIX File Descriptor：

cpp
VmmAllocation allocation =
    create_exportable_allocation(0, requested_size);

int exported_fd = -1;
cuMemExportToShareableHandle(
    &exported_fd,
    allocation.handle,
    CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR,
    0);

File Descriptor 的整数值只在所属进程的 FD Table 中有意义。例如 Process 0 的 fd=17 与 Process 1 自己的 fd=17 没有必然关系，不能使用普通 write(socket, &fd, sizeof(fd)) 传递。Linux 上应通过 Unix Domain Socket 的 SCM_RIGHTS 发送：

cpp
#include <sys/socket.h>

void send_fd(int socket_fd, int fd) {
    char payload = 0;
    struct iovec iov{&payload, sizeof(payload)};

    char control[CMSG_SPACE(sizeof(int))]{};
    struct msghdr message{};
    message.msg_iov = &iov;
    message.msg_iovlen = 1;
    message.msg_control = control;
    message.msg_controllen = sizeof(control);

    struct cmsghdr* cmsg = CMSG_FIRSTHDR(&message);
    cmsg->cmsg_level = SOL_SOCKET;
    cmsg->cmsg_type = SCM_RIGHTS;
    cmsg->cmsg_len = CMSG_LEN(sizeof(int));
    std::memcpy(CMSG_DATA(cmsg), &fd, sizeof(fd));

    sendmsg(socket_fd, &message, 0);
}

int recv_fd(int socket_fd) {
    char payload = 0;
    struct iovec iov{&payload, sizeof(payload)};

    char control[CMSG_SPACE(sizeof(int))]{};
    struct msghdr message{};
    message.msg_iov = &iov;
    message.msg_iovlen = 1;
    message.msg_control = control;
    message.msg_controllen = sizeof(control);

    recvmsg(socket_fd, &message, 0);

    struct cmsghdr* cmsg = CMSG_FIRSTHDR(&message);
    if (cmsg == nullptr ||
        cmsg->cmsg_level != SOL_SOCKET ||
        cmsg->cmsg_type != SCM_RIGHTS) {
        throw std::runtime_error("failed to receive VMM fd");
    }

    int fd = -1;
    std::memcpy(&fd, CMSG_DATA(cmsg), sizeof(fd));
    return fd;
}

SCM_RIGHTS 会在接收进程的 FD Table 中安装一个引用同一内核对象的新 FD。sendmsg 完成后 Exporter 可以关闭自己的 exported_fd，但不能提前释放 CUDA Allocation。

4.3 Import、Map 与 SetAccess ​

Importer 收到 POSIX fd 后，将其恢复为 CUmemGenericAllocationHandle，再映射到自己的 CUDA 虚拟地址空间。Allocation Size 不是 fd 的替代信息，应通过普通控制消息单独传递。

cpp
CUdeviceptr import_vmm_allocation(
    int received_fd,
    size_t allocation_size,
    int consumer_device,
    CUmemGenericAllocationHandle* imported_handle) {
    cuMemImportFromShareableHandle(
        imported_handle,
        reinterpret_cast<void*>(
            static_cast<intptr_t>(received_fd)),
        CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR);

    // Import 成功后不再需要这个 OS Handle。
    close(received_fd);

    CUmemAllocationProp prop{};
    cuMemGetAllocationPropertiesFromHandle(
        &prop,
        *imported_handle);

    size_t granularity = 0;
    cuMemGetAllocationGranularity(
        &granularity,
        &prop,
        CU_MEM_ALLOC_GRANULARITY_MINIMUM);

    CUdeviceptr peer_address = 0;
    cuMemAddressReserve(
        &peer_address,
        allocation_size,
        granularity,
        0,
        0);

    cuMemMap(
        peer_address,
        allocation_size,
        0,
        *imported_handle,
        0);

    set_rw_access(
        peer_address,
        allocation_size,
        consumer_device);

    return peer_address;
}

随后可以把 peer_address 当作普通 Device Pointer 传给 Kernel：

cpp
CUmemGenericAllocationHandle imported_handle;
CUdeviceptr peer_address = import_vmm_allocation(
    received_fd,
    allocation_size,
    1,
    &imported_handle);

size_t count = logical_size / sizeof(float);
add_one<<<(count + 255) / 256, 256>>>(
    reinterpret_cast<float*>(peer_address), count);
cudaDeviceSynchronize();

4.4 清理顺序 ​

Importer 应按以下顺序清理：

cpp
cuMemUnmap(peer_address, allocation_size);
cuMemAddressFree(peer_address, allocation_size);
cuMemRelease(imported_handle);

Exporter 在所有 Importer 完成后清理自己的映射和 Allocation：

cpp
cuMemUnmap(allocation.address, allocation.allocation_size);
cuMemAddressFree(allocation.address, allocation.allocation_size);
cuMemRelease(allocation.handle);

cuMemRelease 释放的是 Allocation Handle 引用。底层物理显存只有在相关 Mapping、Generic Allocation Handle 和导出的 Shareable Handle 都被释放后才会归还给系统。

5. Legacy CUDA IPC 与 VMM-based IPC 的关系 ​

NVIDIA 文档将 cudaIpcGetMemHandle 这一组接口归入 Legacy IPC。VMM-based IPC 提供另一种更显式的 Memory Sharing 模型，两者是并列的公开接口，不是通信算法，也不是上下级调用关系。

两者最直观的工程区别是 API 层级和控制粒度：Legacy CUDA IPC 通常使用较高层的 CUDA Runtime API，例如 cudaIpcGetMemHandle 与 cudaIpcOpenMemHandle，由 CUDA 隐式完成远端 Allocation 的导入和虚拟地址映射；VMM-based IPC 使用更底层的 CUDA Driver API，例如 cuMemCreate、cuMemAddressReserve、cuMemMap 与 cuMemSetAccess，由应用显式管理物理 Allocation、虚拟地址和访问权限。

text
Legacy CUDA IPC / Runtime API
  cudaMalloc
      ↓
  cudaIpcGetMemHandle / cudaIpcOpenMemHandle
      ↓
  Driver 隐式建立 Peer Mapping

VMM-based IPC / Driver API
  cuMemCreate                  Physical Allocation
      +
  cuMemAddressReserve         Virtual Address
      +
  cuMemMap                    VA → Allocation Mapping
      +
  cuMemSetAccess              Per-device Access Permission

Driver API 也提供 cuIpcGetMemHandle、cuIpcOpenMemHandle 等 Legacy CUDA IPC 接口。Legacy CUDA IPC 并非只能通过 Runtime API 使用；它与 VMM-based IPC 的核心区别是前者提供较完整的自动映射抽象，后者通过 Driver API 暴露更细粒度的地址空间控制。

text
跨进程共享 GPU Device Memory
├─ Legacy CUDA IPC
│  ├─ cudaMalloc
│  ├─ cudaIpcGetMemHandle
│  └─ cudaIpcOpenMemHandle
│
└─ VMM-based IPC
   ├─ cuMemCreate
   ├─ cuMemExportToShareableHandle
   ├─ cuMemImportFromShareableHandle
   ├─ cuMemAddressReserve
   ├─ cuMemMap
   └─ cuMemSetAccess
对比项Legacy CUDA IPCVMM-based IPC
常用 API 层级CUDA Runtime API 的 cudaIpc*；也有 Driver API 的 cuIpc*CUDA Driver API 的 cuMem*
映射控制OpenMemHandle 时由 CUDA 隐式完成应用显式执行 Reserve、Map、SetAccess
典型来源cudaMalloc AllocationcuMemCreate Allocation
导出对象完整 CUDA AllocationGeneric Physical Allocation Handle
导出形式cudaIpcMemHandle_tPOSIX fd、Win32 Handle 或 Fabric Handle
VA 控制Driver 自动选择并返回 Importer Pointer应用显式 Reserve、Map、Unmap
访问权限Open 时按需建立 Peer AccesscuMemSetAccess 显式配置
多 Allocation 拼接不直接提供可以把多个 Allocation 映射到连续 VA
使用复杂度较低较高
典型用途固定通信 Buffer、简单单机共享可扩展 Allocator、显式地址布局、跨 API 互操作

同一个 Physical Allocation 不需要先经过 Legacy CUDA IPC 再经过 VMM-based IPC。应用根据 Allocation 的创建方式选择匹配的导出路径：

text
cudaMalloc-backed Pointer
  → cudaIpcGetMemHandle / cudaIpcOpenMemHandle

cuMemCreate + cuMemMap-backed Pointer
  → cuMemExportToShareableHandle
  → cuMemImportFromShareableHandle
  → cuMemMap

cuMemRetainAllocationHandle 可以从已经由 cuMemMap 映射的地址恢复其 Generic Allocation Handle，但它不能把普通 cudaMalloc Pointer 转换成 VMM Allocation。

6. 性能关系 ​

两种机制完成映射后，Kernel 都通过 GPU 虚拟地址访问同一类 Device Memory，远端请求也都经过实际的 NVLink、NVSwitch 或 PCIe 路径。因此，在物理 Allocation、访问权限、地址对齐和 Kernel 完全相同时，稳态 Load/Store 性能通常接近。

主要差异出现在映射建立与完整数据路径：

阶段Legacy CUDA IPCVMM-based IPC
注册Handle 获取与 Open 相对简单Export、OS Handle 传递、Import、Reserve、Map、SetAccess
稳态访问普通 Peer Pointer Load/Store普通 Peer Pointer Load/Store
地址布局Driver 决定应用可控制映射与拼接
碎片与 TLB取决于原始 Allocation还受 Chunk 数量和 Mapping Granularity 影响

公平 Benchmark 应把以下指标分开：

  1. Allocation 与 Handle 注册时间。
  2. 第一次访问或第一次 Kernel 的冷启动时间。
  3. Warmup 后 Kernel-only 延迟和有效带宽。
  4. 是否存在额外 Device-to-Device Staging Copy。
  5. P2P 实际经过 NVLink、NVSwitch 还是 PCIe。

VMM 注册更复杂并不意味着每次 Peer Load 都更慢；同样，某条 VMM 路径更快也可能只是因为它减少了额外 Copy，而不是 cuMemMap 本身提高了链路带宽。

7. 正确性与安全边界 ​

实现共享显存通信时需要同时满足以下约束：

  • 生命周期：Exporter 必须晚于全部 Importer 释放原 Allocation。
  • 生产者—消费者同步：指针可访问不表示数据已经写完，应使用 Event、Stream Dependency、Barrier 或控制消息。
  • Buffer 复用：循环通信需要 Slot 或 Epoch，不能在消费者完成前覆盖数据。
  • 内存顺序：自定义 GPU Flag 协议需要正确的 Release/Acquire 或 System-scope 原子语义，不能只依赖普通 Load/Store 和 volatile 推断跨设备可见性。
  • 错误处理：生产代码必须检查所有 CUDA 返回值，并处理部分注册成功后的回滚。
  • Handle 权限：IPC/VMM Handle 是访问显存的能力凭证，不应发送给不可信进程。
  • 进程模型：避免在已经初始化 CUDA 的多线程进程中直接 fork()；优先使用独立启动的 Worker 或 spawn 模式。

8. 参考 ​

  1. NVIDIA CUDA Programming Guide：Interprocess Communication
  2. NVIDIA CUDA Programming Guide：Virtual Memory Management
  3. NVIDIA CUDA Runtime API：Device Management / IPC
  4. NVIDIA CUDA Driver API：Virtual Memory Management
  5. NVIDIA CUDA Samples：memMapIPCDrv