news 2026/9/16 8:27:44

CUDA simpleIPC 示例深度解析:基于 CUDA Runtime API 的多进程跨 GPU 互操作

作者头像

张小明

前端开发工程师

1.2k 24
文章封面图
CUDA simpleIPC 示例深度解析:基于 CUDA Runtime API 的多进程跨 GPU 互操作

CUDA simpleIPC 示例深度解析:基于 CUDA Runtime API 的多进程跨 GPU 互操作

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples

导读

simpleIPC是 NVIDIA CUDA Samples 仓库(cpp/0_Introduction/simpleIPC)中一个最基础的进程间通信(Inter Process Communication, IPC)示例,它演示了如何让多个操作系统进程各自绑定一块 GPU,并借助 CUDA IPC 机制在进程间共享设备内存指针与 CUDA 事件。通过阅读本文,你将掌握cudaIpcGetMemHandle/cudaIpcOpenMemHandle这一对"导出-导入"设备指针的核心用法、IPC 事件在跨进程流同步中的作用,以及一个可复用的"主进程分发 + 子进程计算 + 共享内存协调"多进程协作架构。

示例概览:一进程一 GPU 的分布式计算雏形

本示例使用 CUDA Runtime API 实现,核心思想是"一个进程对应一块 GPU"(one process per GPU)来完成计算任务,本质上是一个简化版的多进程多 GPU 计算框架:

  • 主进程(parent)枚举系统中所有支持 IPC 的设备,为每块设备分配一段设备内存并创建跨进程事件;
  • 每个子进程(child)绑定一个 GPU,打开其他进程内存的 IPC 句柄,在各自设备上启动同一个simpleKernel内核,交叉写入其他设备的内存;
  • 最后每个子进程把属于自己的缓冲区拷回主机内存进行校验,确认跨进程读写的数据内容正确。

整个流程贯穿了 CUDA Systems Integration、Peer to Peer(P2P)与 InterProcess Communication 三个关键概念(见 README 的 Key Concepts 一节),是理解生产级多进程 CUDA 应用(如多 GPU 推理服务、多进程管线)的入门范本。

支持环境与前置条件

根据 simpleIPC/README.md,本示例的软硬件要求如下:

维度要求
GPU 计算能力Compute Capability 3.0 及以上
操作系统Linux 或 Windows
CPU 架构x86_64、ppc64le(注意:源码中明确对 ARM 返回EXIT_WAIVED,见下文)
构建依赖CUDA Toolkit、IPC 相关支持(见仓库根 README.md 中对 CUDA IPC 的说明)

从 simpleIPC.cu 的main函数可以看到平台限制的落地实现:__arm__/__aarch64__下直接打印 "Not supported on ARM" 并返回EXIT_WAIVED,与 CMakeLists.txt 中 "Will not build sample simpleIPC - not supported on aarch64" 的构建期判断相互印证。

此外,CMakeLists.txt 还展示了本示例对编译工具链的要求:CMake 最低版本 3.20,启用cxx_std_17cuda_std_17,并开启CUDA_SEPARABLE_COMPILATION(可分离编译,便于cudaOccupancyMaxActiveBlocksPerMultiprocessor等运行时 API 查询内核属性)。

核心架构:共享内存驱动的多进程协作

整个程序只有一个二进制文件simpleIPC,通过命令行参数个数区分角色(simpleIPC.cu):

  • 无参数运行:进入parentProcess,扮演主进程;
  • 带一个整数参数运行:进入childProcess(id),扮演编号为id的子进程。

关键数据结构 shmStruct

主进程与子进程之间传递设备选择、IPC 句柄等信息,依赖一块由主进程创建、所有子进程映射的系统共享内存,其结构定义在 simpleIPC.cu:

typedef struct shmStruct_st { size_t nprocesses; // 参与计算的进程/设备数量 int barrier; // 软件屏障计数器 int sense; // 屏障翻转位 int devices[MAX_DEVICES]; // 每个进程绑定的设备号 cudaIpcMemHandle_t memHandle[MAX_DEVICES]; // 每块设备内存的 IPC 句柄 cudaIpcEventHandle_t eventHandle[MAX_DEVICES]; // 每个跨进程事件的 IPC 句柄 } shmStruct;

其中MAX_DEVICES为 32(simpleIPC.cu),DATA_SIZE为 64MB(simpleIPC.cu)。注释特别说明:对直接 NVLINK 或 PCI-E 直连的 GPU,同一时刻最多允许 8 个对等端(peers),而 NVSWITCH 连接(如 DGX-2 类平台)不受此限制。

共享内存与进程管理工具

共享内存的创建/打开/关闭、进程的派生/等待由仓库通用工具 helper_multiprocess.h 与 helper_multiprocess.cpp 提供:

  • Linux 下sharedMemoryCreate/sharedMemoryOpen基于 POSIXshm_open+mmapMAP_SHARED),Windows 下基于CreateFileMapping+MapViewOfFile(helper_multiprocess.cpp);
  • spawnProcess在 Linux 使用fork+execvp,Windows 使用CreateProcess(helper_multiprocess.cpp);
  • 为避免共享内存名冲突,主进程用自身 PID 拼接出唯一的命名(如simpleIPCshm12345),子进程通过getppid()拿到父进程 PID 后打开同一块共享内存(simpleIPC.cu)。

主进程流程:设备筛选、句柄导出与子进程派发

parentProcess(simpleIPC.cu)按以下步骤工作:

1. 设备筛选。遍历所有设备(cudaGetDeviceCount),要求同时满足:

  • 支持统一寻址(prop.unifiedAddressing),这是 CUDA IPC 的前提;
  • 计算模式为cudaComputeModeDefault(不能是独占/禁止模式,因为本示例要求两个进程同时访问同一设备);
  • 与已选设备两两满足cudaDeviceCanAccessPeer(双向)。

满足全部条件才将该设备加入shm->devices数组,否则打印提示并跳过(simpleIPC.cu)。

2. 显式开启 P2P。对入选设备两两调用cudaSetDevice+cudaDeviceEnablePeerAccess。源码注释指出:这一步对 IPC 并非必需,但会在设备上预先建立对等关系,对于"每 GPU 最多 8 个 peer"的系统,相当于提前占用并理顺对等关系(simpleIPC.cu)。

3. 分配内存并导出句柄。在每个入选设备上cudaMalloc一块DATA_SIZE(64MB)内存,然后:

  • cudaIpcGetMemHandle导出设备内存的 IPC 句柄存入共享内存;
  • cudaEventCreate创建事件,标志位为cudaEventDisableTiming | cudaEventInterprocess(跨进程事件必须禁用计时);
  • cudaIpcGetEventHandle导出事件句柄存入共享内存(simpleIPC.cu)。

4. 派发并等待子进程。为每个入选设备spawnProcess一个子进程,参数是设备索引i;随后waitProcess阻塞等待全部子进程退出并检查退出码(simpleIPC.cu)。

子进程流程:句柄导入、跨设备写入与数据校验

childProcess(id)(simpleIPC.cu)是计算与验证的核心:

1. 打开共享内存并绑定设备。打开父进程创建的共享内存,读取进程总数procCount与自己的设备号,然后cudaSetDevice(shm->devices[id]),并创建非阻塞流cudaStreamNonBlocking(simpleIPC.cu)。

2. 导入所有设备的内存与事件句柄。对每个进程i

cudaIpcOpenMemHandle(&ptr, shm->memHandle[i], cudaIpcMemLazyEnablePeerAccess); cudaIpcOpenEventHandle(&event, shm->eventHandle[i]);

这里使用了cudaIpcMemLazyEnablePeerAccess标志——注释明确说明:不需要显式开启 peer access,导入 IPC 句柄时即可惰性建立对等访问(simpleIPC.cu)。这与父进程中的显式cudaDeviceEnablePeerAccess形成对照。

3. 循环跨设备写缓冲。procCount轮循环中,每个子进程按下标(i + id) % procCount选择一块"下一个要写入"的缓冲区:

  • cudaStreamWaitEvent让本流等待该缓冲区当前持有者记录的事件,确保缓冲区就绪;
  • 启动simpleKernel把整块 64MB 缓冲区逐字节写入自身设备编号id
  • cudaEventRecord记录事件,通知后续消费者该缓冲区已写完;
  • barrierWait软屏障同步所有子进程,防止某个进程冲得太快、覆盖掉他人还在等待的事件记录(simpleIPC.cu)。

4. 拷回并校验。等待属于自己的缓冲区事件后,cudaMemcpyAsync把 64MB 数据拷回主机端verification_buffer,逐字节与期望值(id + 1) % procCount比较——因为缓冲区最终写入者正是"下一个编号"的兄弟进程(simpleIPC.cu)。

软件屏障的实现

barrierWait(simpleIPC.cu)是一个两阶段的 CPU 原子屏障:先做"报到"(__sync_add_and_fetch/InterlockedAdd累加计数,最后一人置sense=1),再做"离场"(递减计数,归零时复位sense=0),通过轮询sense实现所有进程的汇合。这是保证事件记录不被兄弟进程抢先覆盖的关键同步点。

CUDA Runtime API 全景

原 README 列出了本示例涉及的全部 CUDA Runtime API,按功能分组整理如下(全部可在 simpleIPC.cu 中找到实际调用):

功能域API
设备管理cudaSetDevicecudaGetDeviceCountcudaGetDevicePropertiescudaDeviceCanAccessPeercudaDeviceEnablePeerAccesscudaOccupancyMaxActiveBlocksPerMultiprocessor
内存管理cudaMalloccudaFreecudaMemcpyAsync
IPC 内存cudaIpcGetMemHandlecudaIpcOpenMemHandlecudaIpcCloseMemHandle
IPC 事件cudaIpcGetEventHandlecudaIpcOpenEventHandle
事件/流cudaEventCreatecudaEventRecordcudaEventSynchronizecudaEventDestroycudaStreamCreateWithFlagscudaStreamWaitEventcudaStreamSynchronizecudaStreamDestroy
错误处理cudaGetLastError

其中"IPC 内存三件套"是贯穿全示例的主线:主进程cudaIpcGetMemHandle导出、子进程cudaIpcOpenMemHandle导入、结束时cudaIpcCloseMemHandle关闭(simpleIPC.cu)。IPC 事件则保证跨进程的异步执行序:事件本身只做同步信号、不记录时间戳,因此创建时必须带cudaEventDisableTiming

构建与运行

本示例的构建配置见 cpp/0_Introduction/simpleIPC/CMakeLists.txt,关键点包括:find_package(CUDAToolkit REQUIRED)、默认 CUDA 架构列表75 80 86 87 89 90 100 110 120、默认开启-lineinfo(调试工具可用的行号信息,ENABLE_CUDA_DEBUG时替换为-G),以及CUDA_SEPARABLE_COMPILATION ON

可参照仓库根 README.md 的通用流程构建(Linux):

mkdir build && cd build cmake .. make -j$(nproc)

构建产物simpleIPC位于 build 目录对应位置,直接无参运行即可:主进程会自动挑选互相可访问的设备、派生子进程,并在所有子进程完成后打印各阶段的 "Step %lld done" 与各进程的 "verifying..." / "Process %d complete!" 输出。若系统没有任何支持 IPC 的设备,程序会输出 "No CUDA devices support IPC" 并以EXIT_WAIVED退出(simpleIPC.cu)。

运行前提与注意事项

  • 统一寻址是硬前提:不支持unifiedAddressing的设备会被直接跳过;
  • 计算模式需为默认模式:独占或禁止模式会使"两进程共享设备"失效;
  • Peer 数量上限:直连拓扑下每 GPU 同时最多 8 个 peer,NVSWITCH 拓扑不设此限;
  • ARM 平台不支持:aarch64 下构建期跳过、运行期EXIT_WAIVED
  • 共享内存命名唯一性:以主进程 PID 为后缀,避免并发运行多个实例时互相干扰。

小结

simpleIPC用不到四百行代码,完整演示了一条"进程间共享设备内存"的可行路径:以操作系统共享内存传递cudaIpcMemHandle_tcudaIpcEventHandle_t,以 IPC 事件 + 软件屏障维持多进程间的数据依赖与执行序,最终通过逐字节校验验证跨进程读写正确性。它同时触及了 CUDA 系统集成(设备枚举、计算模式)、P2P(peer 能力检测与使能)与 IPC 三大主题,是从单进程多流走向多进程多 GPU 编程的关键桥梁。仓库中同类的进阶参考还包括 simpleP2P(P2P 直连)与 simpleMultiGPU(多 GPU 协作),而 helper_multiprocess.cpp 中的共享内存与进程派生封装则可直接复用到你自己的多进程 CUDA 项目中。

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples

创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考

版权声明: 本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若内容造成侵权/违法违规/事实不符,请联系邮箱:809451989@qq.com进行投诉反馈,一经查实,立即删除!
网站建设 2026/9/16 8:25:25

lspci背后:Linux内核PCI子系统与sysfs的完整链路

说实话,在AI Infra这个岗位上待久了,lspci应该是我敲得最频繁的命令之一。新机器到位先敲一遍看GPU/NIC在不在域;驱动装完再敲一遍确认driver绑定;线上掉卡的时候更是要反复敲。但绝大多数时候,我们只看它的输出&#…

作者头像 李华
网站建设 2026/9/16 8:24:59

基于RAG的PHP微信AI客服系统:源码实践与部署解析

最近把一套PHP原创微信AI客服系统的源码完整梳理了一遍,顺手用在了几个客户的项目里,效果超出预期。微信生态里做客服系统并不新鲜,但把AI能力尤其是大模型语义理解融入其中,让机器人真正能扛住全天候的咨询压力,这条路…

作者头像 李华
网站建设 2026/9/16 8:24:57

AI+军事技术:AI在战场的应用与伦理争议

摘要:本文从技术角度系统梳理了AI在军事领域的应用现状与争议。文章首先以2024年以军在加沙冲突中使用的三套AI系统(目标识别、火力规划、态势感知)为切入点,展示了AI如何将"目标到打击"周期从几十分钟缩短到几分钟&…

作者头像 李华
网站建设 2026/9/16 8:24:47

网站在哪里设置关键字避坑指南:5处核心位置与代码实战

网站在哪里设置关键字避坑指南:5处核心位置与代码实战 网站做好了没人访问,是不是你的日常?很多站长盯着服务器日志发呆,发现流量全是0.0,或者只有几个蜘蛛IP。别急着怪算法,先问问自己:你告诉搜索引擎你是谁了吗?…

作者头像 李华
网站建设 2026/9/16 8:24:34

IEA下调250万桶需求,油价为何反而暴涨?

IEA刚把2026年全球石油需求预期下调了250万桶/日。按常理,需求预期走弱,油价该承压。 可市场走出来的结果完全相反:北海Brent从91美元附近一路冲到113.48美元,8月全球库存单月砍掉约9500万桶,柴油裂差甚至创下纪录。 真…

作者头像 李华