企业官网建设流程全解析

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 Management，VMM）相关的cuMemMap系列 API，在一个进程对应一块 GPU（one 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 或 Windows（README 的 Supported OSes 一栏同时列出 Linux、Windows、QNX）。

关键实现文件位于仓库的 cpp/3_CUDA_Features/memMapIPCDrv 目录：

支持范围与环境依赖

支持的 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、QNX；
• CPU 架构：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 额外链接rt（POSIX 共享内存shm_open所需），QNX 额外链接socket；
• 使用add_custom_command调用nvcc -ptx将 memMapIpc_kernel.cu 预编译为memMapIpc_kernel64.ptx，运行时由驱动加载（详见后文模块加载小节）；
• 源码中通过宏选择 PTX 文件：64 位平台使用memMapIpc_kernel64.ptx，32 位平台使用memMapIpc_kernel32.ptx（见 memMapIpc.cpp）。

整体架构：父进程分配导出，子进程导入计算

程序入口 main 首先调用cuInit(0)初始化驱动，然后根据命令行参数分派：

• 无参数运行时，进入parentProcess(argv[0])：扮演"协调者"角色；
• 带两个参数（设备号、进程序号）运行时，进入childProcess(devId, id, argv)：扮演"计算者"角色。父进程通过spawnProcess以fork + exec（POSIX）或CreateProcess（Windows）方式拉起子进程，见 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）

1. 设备枚举与筛选（memMapIpc.cpp）：
   • cuDeviceGetCount获取设备总数；
   • 逐设备检查CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED（必须支持虚拟地址管理）、计算模式必须为CU_COMPUTEMODE_DEFAULT（独占/禁止模式不符合本样例多进程共享前提）、以及对应平台 IPC 句柄类型属性（Linux/QNX 检查CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR_SUPPORTED，Windows 检查CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_HANDLE_SUPPORTED）；
   • 通过cuDeviceCanAccessPeer双向检查与已选设备的 peer 能力，保证选出的设备集合两两可互访；
   • 为每个入选设备创建上下文并用cuCtxEnablePeerAccess建立双向 peer 访问。源码注释特别说明：这一步对 IPC 本身不是必需的，但有助于清理 peer 关系，规避单 GPU 最多 8 个同时 peer 的硬件限制（NVSwitch 互联如 DGX-2 不受此限制）。
2. 分配与导出：调用memMapAllocateAndExportMemory，在"首选设备"（selectedDevices[0]）上为每个进程分配一块 4MB 物理内存并导出为可共享句柄。
3. 创建子进程并分发句柄：为每个入选设备拉起一个子进程，通过共享内存中的屏障（barrier）同步就绪状态，然后经ipcCreateSocket+ipcSendShareableHandles把全部句柄分发给每个子进程。
4. 回收：等待所有子进程退出、cuMemRelease释放句柄、关闭 socket 与共享内存。

子进程流程（childProcess）

1. 打开共享内存读取nprocesses，与父进程在屏障处汇合；
2. 通过ipcRecvShareableHandles接收父进程传来的全部可共享句柄；
3. cuDeviceGet+cuCtxCreate创建自己的设备上下文和流；
4. 加载内核模块、按占用率计算网格规模；
5. cuMemAddressReserve预留连续虚拟地址空间；
6. memMapImportAndMapMemory导入句柄并映射、设置访问权限；
7. 逐块启动内核写数据，每步之间用屏障同步以保证数据确定性；
8. cuMemcpyDtoHAsync拷回主机校验内容；
9. 依次执行cuStreamDestroy、cuModuleUnload、cuCtxDestroy、memMapUnmapAndFreeMemory完成清理。

核心一：分配与导出 —— cuMemCreate + cuMemExportToShareableHandle

物理内存的分配与导出集中在memMapAllocateAndExportMemory（memMapIpc.cpp）：

CUmemAllocationProp 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.cpp）：Linux/QNX 使用CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR（文件描述符），Windows 使用CU_MEM_HANDLE_TYPE_WIN32（NT HANDLE）；
• 粒度检查：cuMemCreate的分配尺寸必须是设备最小分配粒度的整数倍，示例数据块 4MB 通常满足该约束；源码在非整数倍时直接打印错误并退出；
• Windows 安全描述符：当使用CU_MEM_HANDLE_TYPE_WIN32时，getDefaultSecurityDescriptor（memMapIpc.cpp）通过ConvertStringSecurityDescriptorToSecurityDescriptorA构造LPSECURITYATTRIBUTES并挂到prop.win32HandleMetaData上，用以限定导出分配允许跨进程传递的范围；其他句柄类型传 NULL 即可。

核心二：导入与映射 —— cuMemImportFromShareableHandle + cuMemMap + cuMemSetAccess

子进程侧的memMapImportAndMapMemory（memMapIpc.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));

memMapUnmapAndFreeMemory（memMapIpc.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与ipcRecvShareableHandle，helper_multiprocess.cpp）；socket 名称默认放在系统临时目录（QNX 因 SDP 8.0.3 起限制，固定使用/storage，见 helper_multiprocess.h）；
• Windows：基于Mailslot（邮槽）通道，父进程用DuplicateHandle把句柄复制进目标进程的地址空间后经WriteFile/ReadFile传递（见 helper_multiprocess.cpp）；
• 共享内存同步区：父进程用shm_open + mmap（POSIX）或CreateFileMapping + MapViewOfFile（Windows）创建一块名为memmap_ipc_shm<pid>的共享内存，存放进程数与屏障计数（shmStruct），见 helper_multiprocess.cpp。

父进程在ipcSendShareableHandles中会把全部句柄广播给每一个子进程（helper_multiprocess.cpp），因此每个子进程都能映射到所有共享缓冲区；子进程完成导入后立即ipcCloseShareableHandle关闭句柄，只保留 VA 映射。

核心五：多进程栅栏同步与数据确定性

为了保证"每块缓冲区的内容由哪个进程写入"是可预期的，示例在 memMapIpc.cpp 实现了基于原子操作的栅栏barrierWait：

static 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-out）sense-reversing 屏障：cpu_atomic_add32在 Linux/QNX 上展开为 GCC 内建原子__sync_add_and_fetch，在 Windows 上展开为InterlockedAdd（memMapIpc.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写入整块缓冲区，配合cuLaunchKernel+cuStreamSynchronize执行；网格规模由cuOccupancyMaxActiveBlocksPerMultiprocessor乘以上下文的多处理器数得出。

内核模块的加载方式：PTX 运行时 JIT

示例通过memMapGetDeviceFunction（memMapIpc.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、cuMemSetAccess
• IPC 句柄：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