cuda-samples memMapIPCDrv 深度解析：基于 cuMemMap 与 Driver API 的多进程跨 GPU 内存共享实战

cuda-samples memMapIPCDrv 深度解析基于 cuMemMap 与 Driver API 的多进程跨 GPU 内存共享实战【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读memMapIPCDrv 是 cuda-samples 仓库中一个极具代表性的 CUDA Driver API 示例它展示了如何借助虚拟内存管理Virtual Memory ManagementVMM相关的cuMemMap系列 API在一个进程对应一块 GPUone process per GPU的架构下实现进程间通信IPC。读完本文你将掌握从物理内存分配、可共享句柄导出、跨进程传输与导入到虚拟地址预留、映射与访问控制、乃至多进程栅栏同步的完整技术链路并能在 Linux / Windows / QNX 上独立构建与运行该示例进行验证。示例定位与核心概念依据 memMapIPCDrv 的 README该示例是一个非常基础very basic的 Driver API 演示程序核心要点如下实现方式使用cuMemMap系列 API 完成进程间通信每个 GPU 由一个独立进程负责计算接口形态完全基于CUDA Driver API而非 Runtime API源码中直接调用cuMemCreate、cuMemMap、cuMemExportToShareableHandle等底层接口核心概念CUDA Driver API、cuMemMap IPC、MMAP内存映射硬件要求Compute Capability 3.0 或更高操作系统Linux 或 WindowsREADME 的 Supported OSes 一栏同时列出 Linux、Windows、QNX。关键实现文件位于仓库的 cpp/3_CUDA_Features/memMapIPCDrv 目录文件职责memMapIpc.cpp主程序父进程负责设备枚举、分配与导出内存、拉起子进程子进程负责导入映射并执行内核memMapIpc_kernel.cu设备端内核向共享缓冲区写入指定值用于验证跨进程数据可达CMakeLists.txt构建脚本将内核编译为 PTX 并链接驱动库helper_multiprocess.h / helper_multiprocess.cpp跨平台多进程与可共享句柄传输工具库支持范围与环境依赖支持的 SM 架构README 明确列出的受支持 SM 架构为SM 5.0、5.2、5.3、6.0、6.1、7.0、7.2、7.5、8.0、8.6、8.7、8.9、9.0。与之对应memMapIPCDrv/CMakeLists.txt 中声明的默认 CUDA 架构集合为75 80 86 87 89 90 100 110 120构建时会针对这些架构生成 PTX/二进制实际运行由驱动按需编译适配。支持的操作系统与 CPU 架构操作系统Linux、Windows、QNXCPU 架构x86_64、armv7l、aarch64。功能依赖与前置条件README 将本示例归类到仓库根 README 的CUDA Interprocess Communication能力项下对应根 README 中的 CUDA Interprocess Communication 一节该节对 IPC 的定位是IPC (Interprocess Communication) allows processes to share device pointers即通过 IPC 让多个进程共享设备指针所指向的内存。前置条件为按平台下载并安装对应的 CUDA Toolkit该链接为 NVIDIA 官方下载页本文仅转述 README 原意不再展开并确保上述 IPC 相关依赖就绪。构建与运行该示例使用 CMake 构建遵循仓库根 README 的 Linux 构建流程# 从仓库根目录或任意样例子目录执行 mkdir build cd build cmake .. # 要求 CMake 3.20 或更高版本 make -j$(nproc) # 编译全部样例构建完成后可执行文件memMapIPCDrv位于build/下对应目录中直接运行即可。Windows 平台则可在x64 Native Tools Command Prompt for VS中执行cmake .. -G Visual Studio 16 2019 -A x64生成解决方案后编译。CMakeLists.txt 还揭示了几个值得注意的构建细节可执行文件由memMapIpc.cpp与../../../Common/helper_multiprocess.cpp共同编译链接而成通过CUDA::cuda_driver链接驱动库Linux 额外链接rtPOSIX 共享内存shm_open所需QNX 额外链接socket使用add_custom_command调用nvcc -ptx将 memMapIpc_kernel.cu 预编译为memMapIpc_kernel64.ptx运行时由驱动加载详见后文模块加载小节源码中通过宏选择 PTX 文件64 位平台使用memMapIpc_kernel64.ptx32 位平台使用memMapIpc_kernel32.ptx见 memMapIpc.cpp。整体架构父进程分配导出子进程导入计算程序入口 main 首先调用cuInit(0)初始化驱动然后根据命令行参数分派无参数运行时进入parentProcess(argv[0])扮演协调者角色带两个参数设备号、进程序号运行时进入childProcess(devId, id, argv)扮演计算者角色。父进程通过spawnProcess以fork execPOSIX或CreateProcessWindows方式拉起子进程见 helper_multiprocess.cpp。关键常量定义在 memMapIpc.cpp#define MAX_DEVICES (32) // 最多参与的设备数NVLink/PCIe 直连 peer 上限 #define PROCESSES_PER_DEVICE 1 // 每个 GPU 一个进程 #define DATA_BUF_SIZE 4ULL * 1024ULL * 1024ULL // 每个共享缓冲区 4MB static const char ipcName[] memmap_ipc_pipe; // 句柄传输通道名 static const char shmName[] memmap_ipc_shm; // 共享内存名会拼接父进程 PID 保证唯一PROCESSES_PER_DEVICE被定义为 1即一卡一进程模型最终选定的设备数为nprocesses受MAX_DEVICES限制。父进程流程parentProcess设备枚举与筛选memMapIpc.cppcuDeviceGetCount获取设备总数逐设备检查CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED必须支持虚拟地址管理、计算模式必须为CU_COMPUTEMODE_DEFAULT独占/禁止模式不符合本样例多进程共享前提、以及对应平台 IPC 句柄类型属性Linux/QNX 检查CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR_SUPPORTEDWindows 检查CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_HANDLE_SUPPORTED通过cuDeviceCanAccessPeer双向检查与已选设备的 peer 能力保证选出的设备集合两两可互访为每个入选设备创建上下文并用cuCtxEnablePeerAccess建立双向 peer 访问。源码注释特别说明这一步对 IPC 本身不是必需的但有助于清理 peer 关系规避单 GPU 最多 8 个同时 peer 的硬件限制NVSwitch 互联如 DGX-2 不受此限制。分配与导出调用memMapAllocateAndExportMemory在首选设备selectedDevices[0]上为每个进程分配一块 4MB 物理内存并导出为可共享句柄。创建子进程并分发句柄为每个入选设备拉起一个子进程通过共享内存中的屏障barrier同步就绪状态然后经ipcCreateSocketipcSendShareableHandles把全部句柄分发给每个子进程。回收等待所有子进程退出、cuMemRelease释放句柄、关闭 socket 与共享内存。子进程流程childProcess打开共享内存读取nprocesses与父进程在屏障处汇合通过ipcRecvShareableHandles接收父进程传来的全部可共享句柄cuDeviceGetcuCtxCreate创建自己的设备上下文和流加载内核模块、按占用率计算网格规模cuMemAddressReserve预留连续虚拟地址空间memMapImportAndMapMemory导入句柄并映射、设置访问权限逐块启动内核写数据每步之间用屏障同步以保证数据确定性cuMemcpyDtoHAsync拷回主机校验内容依次执行cuStreamDestroy、cuModuleUnload、cuCtxDestroy、memMapUnmapAndFreeMemory完成清理。核心一分配与导出 —— cuMemCreate cuMemExportToShareableHandle物理内存的分配与导出集中在memMapAllocateAndExportMemorymemMapIpc.cppCUmemAllocationProp prop {}; prop.type CU_MEM_ALLOCATION_TYPE_PINNED; // 设备端固定的物理分配 prop.location.type CU_MEM_LOCATION_TYPE_DEVICE; prop.location.id (int)backingDevice; // 后备物理设备 prop.requestedHandleTypes ipcHandleTypeFlag; // 声明可导出句柄类型 // 查询该设备支持的最小分配粒度 checkCudaErrors(cuMemGetAllocationGranularity(granularity, prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM)); if (allocSize % granularity) { /* 非粒度倍数则退出 */ } // 逐块创建分配并导出 checkCudaErrors(cuMemCreate(allocationHandles[i], allocSize, prop, 0)); checkCudaErrors(cuMemExportToShareableHandle((void *)shareableHandles[i], allocationHandles[i], ipcHandleTypeFlag, 0));几个要点句柄类型按平台选择memMapIpc.cppLinux/QNX 使用CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR文件描述符Windows 使用CU_MEM_HANDLE_TYPE_WIN32NT HANDLE粒度检查cuMemCreate的分配尺寸必须是设备最小分配粒度的整数倍示例数据块 4MB 通常满足该约束源码在非整数倍时直接打印错误并退出Windows 安全描述符当使用CU_MEM_HANDLE_TYPE_WIN32时getDefaultSecurityDescriptormemMapIpc.cpp通过ConvertStringSecurityDescriptorToSecurityDescriptorA构造LPSECURITYATTRIBUTES并挂到prop.win32HandleMetaData上用以限定导出分配允许跨进程传递的范围其他句柄类型传 NULL 即可。核心二导入与映射 —— cuMemImportFromShareableHandle cuMemMap cuMemSetAccess子进程侧的memMapImportAndMapMemorymemMapIpc.cpp完成句柄 → CUDA 句柄 → 虚拟地址 → 访问权限的四步转换CUmemAccessDesc accessDescriptor; accessDescriptor.location.type CU_MEM_LOCATION_TYPE_DEVICE; accessDescriptor.location.id mapDevice; // 映射目标设备 accessDescriptor.flags CU_MEM_ACCESS_FLAGS_PROT_READWRITE; // 读写权限 for (int i 0; i shareableHandles.size(); i) { // 1. 从平台句柄导入为 CUDA 分配句柄 checkCudaErrors(cuMemImportFromShareableHandle(allocationHandles[i], (void *)(uintptr_t)shareableHandles[i], ipcHandleTypeFlag)); // 2. 映射到预留的虚拟地址区间d_ptr i * mapSize checkCudaErrors(cuMemMap(d_ptr (i * mapSize), mapSize, 0, allocationHandles[i], 0)); // 3. 映射完成后即可释放分配句柄后备内存由映射维系 checkCudaErrors(cuMemRelease(allocationHandles[i])); } // 4. 为整个 VA 区间设置访问描述符含 peer 访问 checkCudaErrors(cuMemSetAccess(d_ptr, shareableHandles.size() * mapSize, accessDescriptor, 1));这里体现了 VMM 模型下分配allocation与映射mapping分离的设计cuMemImportFromShareableHandle得到的是分配句柄cuMemMap才把物理后备存储与虚拟地址绑定cuMemRelease释放分配句柄后只要映射仍然存在内存就不会被回收源码注释明确指出The allocation will be kept live until it is unmapped。最后统一cuMemSetAccess把整段 VA 区间以读写权限映射到目标设备实现对 peer 设备的访问开放。核心三虚拟地址生命周期管理子进程在导入之前先预留 VA 空间退出前再统一回收// 预留连续 VA按 4MB 对齐预留 nprocesses * 4MB checkCudaErrors(cuMemAddressReserve(d_ptr, procCount * DATA_BUF_SIZE, DATA_BUF_SIZE, 0, 0));memMapUnmapAndFreeMemorymemMapIpc.cpp则执行对称的清理// 解除映射由于句柄已 cuMemRelease且这是最后一处引用后备存储随之释放 checkCudaErrors(cuMemUnmap(dptr, size)); // 释放 VA 区间允许未来 cuMemAddressReserve 或 malloc/mmap 复用该地址 checkCudaErrors(cuMemAddressFree(dptr, size));源码注释还提醒解除映射后若继续访问该 VA 区间将触发 fault直到重新映射。这正对应 README 所列 API 清单中的cuMemAddressReserve、cuMemUnmap、cuMemAddressFree三个接口完整展示了虚拟内存预留 → 使用 → 释放的闭环。核心四可共享句柄的跨进程传输机制句柄传输不是 CUDA API 的职责而是由示例自带的 helper_multiprocess.cpp 实现的跨平台机制这使该示例具备极高的教学价值Linux / QNX基于AF_UNIX SOCK_DGRAM数据报 socket利用sendmsg/recvmsg的SCM_RIGHTS附属数据ancillary data把文件描述符fd原样传送给目标进程见ipcSendShareableHandle与ipcRecvShareableHandlehelper_multiprocess.cppsocket 名称默认放在系统临时目录QNX 因 SDP 8.0.3 起限制固定使用/storage见 helper_multiprocess.hWindows基于Mailslot邮槽通道父进程用DuplicateHandle把句柄复制进目标进程的地址空间后经WriteFile/ReadFile传递见 helper_multiprocess.cpp共享内存同步区父进程用shm_open mmapPOSIX或CreateFileMapping MapViewOfFileWindows创建一块名为memmap_ipc_shmpid的共享内存存放进程数与屏障计数shmStruct见 helper_multiprocess.cpp。父进程在ipcSendShareableHandles中会把全部句柄广播给每一个子进程helper_multiprocess.cpp因此每个子进程都能映射到所有共享缓冲区子进程完成导入后立即ipcCloseShareableHandle关闭句柄只保留 VA 映射。核心五多进程栅栏同步与数据确定性为了保证每块缓冲区的内容由哪个进程写入是可预期的示例在 memMapIpc.cpp 实现了基于原子操作的栅栏barrierWaitstatic void barrierWait(volatile int *barrier, volatile int *sense, unsigned int n) { int count cpu_atomic_add32(barrier, 1); // 登记进场 if (count n) *sense 1; // 最后一人进场放行 while (!*sense) ; count cpu_atomic_add32(barrier, -1); // 登记离场 if (count 0) *sense 0; // 最后一人离场复位 while (*sense) ; }这是一个经典的两阶段check-in / check-outsense-reversing 屏障cpu_atomic_add32在 Linux/QNX 上展开为 GCC 内建原子__sync_add_and_fetch在 Windows 上展开为InterlockedAddmemMapIpc.cpp。子进程在计算循环中每个进程写入的缓冲区编号按(i id) % procCount轮转memMapIpc.cpp每写完一块都等待全体进程到位后再进入下一步从而保证最终每块缓冲区的值都来自编号紧邻其后的兄弟进程校验阶段据此验证// 期望值 紧邻的下一个进程 id char compareId (char)((id 1) % procCount); for (unsigned long long j 0; j DATA_BUF_SIZE; j) { if (verification_buffer[j] ! compareId) { /* 打印不匹配位置并中止 */ } }设备端内核 memMapIpc_kernel.cu 只是一个填充内核按blockIdx.x * blockDim.x threadIdx.x的线性索引以 stride 方式把val写入整块缓冲区配合cuLaunchKernelcuStreamSynchronize执行网格规模由cuOccupancyMaxActiveBlocksPerMultiprocessor乘以上下文的多处理器数得出。内核模块的加载方式PTX 运行时 JIT示例通过memMapGetDeviceFunctionmemMapIpc.cpp加载内核体现了 Driver API 特有的模块管理能力先用findModulePath在可执行文件所在路径查找memMapIpc_kernel64.ptx或 32 位变体找不到则回退到memMapIpc_kernel.cubin找到PTX时走cuModuleLoadDataEx并配置 3 个 JIT 选项CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES日志缓冲大小 1024 字节、CU_JIT_INFO_LOG_BUFFER日志缓冲指针、CU_JIT_MAX_REGISTERS最大寄存器数 32加载后打印 PTX JIT 日志找到CUBIN时走cuModuleLoad直接加载二进制最后用cuModuleGetFunction取得memMapIpc_kernel的函数句柄供cuLaunchKernel使用。这一分支与 CMake 中把内核提前编译为 PTX的自定义命令相衔接构建期只产出 PTX运行时由驱动针对实际 GPU 做 JIT 编译这正是 Driver API 模式下典型的部署形态。涉及的 CUDA Driver API 一览README 完整罗列了本示例用到的 Driver API可按需查阅对应驱动头文件与文档初始化/枚举cuInit、cuDeviceGet、cuDeviceGetCount、cuDeviceGetAttribute、cuDeviceCanAccessPeer上下文与流cuCtxCreate、cuCtxDestroy、cuCtxSetCurrent、cuCtxEnablePeerAccess、cuStreamCreate、cuStreamDestroy、cuStreamSynchronize虚拟内存管理cuMemCreate、cuMemRelease、cuMemGetAllocationGranularity、cuMemAddressReserve、cuMemAddressFree、cuMemMap、cuMemUnmap、cuMemSetAccessIPC 句柄cuMemExportToShareableHandle、cuMemImportFromShareableHandle模块与内核cuModuleLoad、cuModuleLoadDataEx、cuModuleGetFunction、cuLaunchKernel、cuOccupancyMaxActiveBlocksPerMultiprocessor数据搬运cuMemcpyDtoHAsync小结与延伸memMapIPCDrv 以极小的代码规模完整演绎了 CUDA 虚拟内存管理与 IPC 的全部关键环节父进程负责cuMemCreate分配并cuMemExportToShareableHandle导出句柄子进程通过cuMemImportFromShareableHandle导入、cuMemAddressReserve/cuMemMap映射、cuMemSetAccess授权最终由cuMemUnmap/cuMemAddressFree释放。结合helper_multiprocess提供的 fd/句柄传输与共享内存屏障读者可以据此在自己的多进程、多 GPU 应用中落地一卡一进程、共享物理后备存储的架构。仓库中同目录的其他样例如simpleIPC、streamOrderedAllocationIPC、memMapIPCDrv的 Python 对应实现 python/4_DistributedComputing/ipcMemoryPool可作为对照进一步理解 Runtime API 与 Driver API 两种路径下 IPC 用法的异同。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考