cuda-samples vectorAddMMAP 深度解析：用 cuMemMap 虚拟内存管理实现向量加法
【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples
导读
vectorAddMMAP是 NVIDIA cuda-samples 仓库中位于 cpp/0_Introduction/vectorAddMMAP 目录下的入门级示例，其核心目标是将 vectorAddDrv 中的传统设备内存分配（cuMemAlloc）替换为基于cuMemMap的虚拟内存映射分配。通过本文，你将掌握cuMemMap/cuMemCreate/cuMemSetAccess等 CUDA Driver API 虚拟内存管理接口的完整调用链，理解"物理内存属性与虚拟地址连续性解耦"的设计思想，并学会如何用虚拟地址区间（VA）组织跨设备驻留、多设备可访问的内存分配，同时保持程序原有访问结构不变。
示例概览：从 cuMemAlloc 到 cuMemMap
原 README 明确指出：该示例用 cuMemMap 分配的设备内存替换了 vectorAddDrv 示例中的设备分配。其核心意义在于——cuMemMapAPI 允许用户在保留内存访问的连续性的同时，自由指定内存的物理属性（例如内存驻留在哪个设备、以何种粒度分配），因此不需要改变程序原有的结构。
对比两个示例的主机端代码可以更直观地看到差异：
- vectorAddDrv.cpp 使用
cuMemAlloc(&d_A, size)三行分配三个设备缓冲区； - vectorAddMMAP.cpp 改用
simpleMallocMultiDeviceMmap(&d_A, &allocationSize, size, backingDevices, mappingDevices)进行映射式分配。
两者在分配完成之后的代码几乎完全一致（cuMemcpyHtoD拷贝输入、cuLaunchKernel启动内核、cuMemcpyDtoH取回结果），这恰好印证了 README 的描述：切换到cuMemMap不改变程序结构，只替换"分配"这一环节。
前置条件与支持平台
- 前置条件：下载并安装对应平台的 CUDA Toolkit（官方提供的安装入口，仓库内代码依赖其驱动与头文件）。
- 支持的 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。
- 支持的操作系统：Linux、Windows。
- 支持的 CPU 架构：x86_64、ppc64le。
- 特殊限制：该示例不支持 aarch64。在 CMakeLists.txt 中，当
CMAKE_SYSTEM_PROCESSOR为aarch64时会打印 "Will not build sample vectorAddMMAP - not supported on aarch64" 并跳过构建。
此外，运行时代码还会检查设备属性CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED，只有支持虚拟地址管理的设备才能使用这些 API（详见下文源码解析）。
涉及的核心 CUDA Driver API
README 列出了该示例使用的全部 Driver API，这里给出每个接口在本示例中的职责：
| API | 在示例中的用途 |
|---|---|
cuInit | 初始化 CUDA Driver API 运行环境 |
cuDeviceGetCount | 统计系统中 GPU 数量，用于收集可作为后备（backing）设备的候选 |
cuDeviceGetAttribute | 查询CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED，判断设备是否支持虚拟地址管理 |
cuDeviceCanAccessPeer | 判断其他设备能否与当前设备建立 P2P 互访，决定哪些设备可作为后备设备 |
cuCtxCreate | 在目标设备上创建 CUDA 上下文 |
cuCtxDestroy | 释放上下文 |
cuMemGetAllocationGranularity | 查询每个参与设备的最小分配粒度，取最大值作为统一粒度 |
cuMemAddressReserve | 预留一段连续的虚拟地址空间（VA 区间） |
cuMemCreate | 按指定物理属性（Pinned + 设备位置）创建物理内存分配句柄 |
cuMemMap | 将物理分配映射到预留的 VA 区间 |
cuMemSetAccess | 为映射设备设置读写访问权限，实现跨设备可见性 |
cuMemRelease | 释放分配句柄（映射建立后句柄不再需要） |
cuMemUnmap | 解除 VA 区间上的映射，释放物理后备存储 |
cuMemAddressFree | 释放 VA 区间，使其可被复用 |
cuModuleLoadData | 从 fatbin 二进制数据加载 CUDA 模块 |
cuModuleGetFunction | 从模块中获取内核函数句柄VecAdd_kernel |
cuLaunchKernel | 启动向量加法内核 |
cuMemcpyHtoD/cuMemcpyDtoH | 主机与设备之间的数据拷贝 |
主机端主流程解析
vectorAddMMAP.cpp 的主函数流程如下：
- 初始化与设备选择：
cuInit(0)初始化驱动，findCudaDeviceDRV（定义于 Common/helper_cuda_drvapi.h）根据命令行-device=参数或按最高 Gflops 自动选择设备。 - 虚拟地址管理能力检查：
checkCudaErrors( cuDeviceGetAttribute(&attributeVal, CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED, cuDevice)); printf("Device %d VIRTUAL ADDRESS MANAGEMENT SUPPORTED = %d.\n", cuDevice, attributeVal); if (attributeVal == 0) { printf("Device %d doesn't support VIRTUAL ADDRESS MANAGEMENT.\n", cuDevice); exit(EXIT_WAIVED); }若设备不支持虚拟地址管理，程序以EXIT_WAIVED（值为 2，定义于 Common/helper_cuda_drvapi.h）退出，表示跳过而非失败。
- 收集后备设备：
getBackingDevices(cuDevice)（vectorAddMMAP.cpp）通过cuDeviceGetCount遍历所有设备，用cuDeviceCanAccessPeer筛出与当前设备支持 P2P 互访的设备，再用虚拟地址管理属性做二次过滤，最终得到可为其分配物理内存的设备列表。 - 创建上下文与加载模块：
cuCtxCreate创建上下文；findFatbinPath定位构建期生成的vectorAdd_kernel64.fatbin（默认宏FATBIN_FILE）；cuModuleLoadData加载二进制模块，cuModuleGetFunction获取VecAdd_kernel函数句柄。 - 分配设备内存：三次调用
simpleMallocMultiDeviceMmap为d_A、d_B、d_C分配虚拟连续、可被mappingDevices读写的内存。 - 数据搬运与内核启动：
cuMemcpyHtoD拷贝输入；以threadsPerBlock = 256、blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock的网格配置，通过参数数组void *args[] = {&d_A, &d_B, &d_C, &N}调用cuLaunchKernel。 - 结果校验：
cuMemcpyDtoH取回h_C，逐元素校验fabs(h_C[i] - (h_A[i] + h_B[i])) > 1e-7f，全部通过则输出Result = PASS。 - 清理：
CleanupNoFailure依次simpleFreeMultiDeviceMmap释放设备内存、free主机内存、cuModuleUnload卸载模块、cuCtxDestroy销毁上下文。
多设备 mmap 分配实现：simpleMallocMultiDeviceMmap
核心分配逻辑封装在 multidevicealloc_memmap.cpp 的simpleMallocMultiDeviceMmap中，接口声明见 multidevicealloc_memmap.hpp。函数参数含义：
dptr（输出）：预留的虚拟地址起始值；allocationSize（输出）：实际预留的 VA 空间大小，释放时必须传回；size（输入）：期望的最小分配字节数，会向上取整；residentDevices（输入）：物理内存需要跨哪些设备条带化驻留（stripe）；mappingDevices（输入）：哪些设备需要读写这块内存；align（输入，默认 0）：额外对齐要求。
虚拟地址布局
头文件注释给出了 VA 映射的可视化布局：
v-stripeSize-v v-rounding -v +-----------------------------------------+ | D1 | D2 | D3 | +-----------------------------------------+ ^-- dptr ^-- dptr + size每个residentDevices中的设备获得等长的条带（stripe），末尾多余空间用于满足所有设备的最小粒度要求。
实现步骤与关键点
- 构造分配属性：
CUmemAllocationProp prop设置type = CU_MEM_ALLOCATION_TYPE_PINNED、location.type = CU_MEM_LOCATION_TYPE_DEVICE，即创建设备锁页（pinned）内存。 - 计算统一最小粒度：对所有
residentDevices与mappingDevices调用cuMemGetAllocationGranularity(..., CU_MEM_ALLOC_GRANULARITY_MINIMUM)，取各设备粒度的最大值作为min_granularity。 - 向上取整：
size = round_up(size, residentDevices.size() * min_granularity)，使总大小可被设备数整除且每个条带都满足粒度；stripeSize = size / residentDevices.size()；通过allocationSize把取整后的大小回传给调用方（释放时使用）。 - 预留 VA 区间：
cuMemAddressReserve(dptr, size, align, 0, 0)预留连续虚拟地址。 - 逐设备创建并映射：循环中对每个
residentDevices[idx]设置prop.location.id，cuMemCreate(&allocationHandle, stripeSize, &prop, 0)创建物理分配，cuMemMap(*dptr + stripeSize * idx, stripeSize, 0, allocationHandle, 0)映射到对应 VA 偏移，随后立即cuMemRelease(allocationHandle)——映射建立后句柄不再需要，物理内存由映射保持存活。 - 设置跨设备访问权限：为每个
mappingDevices构造CUmemAccessDesc（location.type = CU_MEM_LOCATION_TYPE_DEVICE、flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE），一次cuMemSetAccess(*dptr, size, descriptors, count)为整个 VA 区间授予读写权限。 - 失败回滚：任一步失败跳转
done标签，若*dptr非空则调用simpleFreeMultiDeviceMmap清理已建立的映射。
值得注意的源码注释（vectorAddMMAP.cpp）强调：即使后备设备与映射设备不同，也无需调用cuCtxEnablePeerAccess，因为cuMemSetAccess显式指定了跨设备映射，但该调用仍受cuDeviceCanAccessPeer约束（这正是先前收集backingDevices时先做 P2P 检查的原因）。
释放实现：simpleFreeMultiDeviceMmap
multidevicealloc_memmap.cpp 中的释放分两步：
cuMemUnmap(dptr, size)：解除整个 VA 区间上的映射。由于句柄此前已通过cuMemRelease释放且这是唯一引用该后备存储的映射，后备物理内存在此处被自动释放；此后访问该 VA 区间将触发错误（fault）。cuMemAddressFree(dptr, size)：归还虚拟地址区间，使其可被后续cuMemAddressReserve或其他操作系统级分配（如malloc、mmap）复用。
设备端内核
内核定义在 vectorAdd_kernel.cu，是经典的逐元素向量加法：
extern "C" __global__ void VecAdd_kernel(const float *A, const float *B, float *C, int N) { int i = blockDim.x * blockIdx.x + threadIdx.x; if (i < N) C[i] = A[i] + B[i]; }extern "C"保证函数名不被名字改编（name mangling），从而可被主机端cuModuleGetFunction(&vecAdd_kernel, cuModule, "VecAdd_kernel")精确查找到。
构建与运行
构建配置位于 CMakeLists.txt：
- 要求 CMake 3.20+，启用
C CXX CUDA三种语言，通过find_package(CUDAToolkit REQUIRED)定位 CUDA 工具包。 - 默认架构列表为
75 80 86 87 89 90 100 110 120，并通过add_custom_command调用nvcc ... -fatbin把 vectorAdd_kernel.cu 编译为vectorAdd_kernel64.fatbin（fatbin 文件生成后由findFatbinPath在运行时定位加载）。 - 可执行文件由
vectorAddMMAP.cpp与multidevicealloc_memmap.cpp组成，链接CUDA::cuda_driver。 - 该目录通过 cpp/0_Introduction/CMakeLists.txt 中的
add_subdirectory(vectorAddMMAP)接入整个仓库的构建体系。
典型构建与运行方式（在仓库根目录）：
mkdir build && cd build cmake .. -DCMAKE_BUILD_TYPE=Release make vectorAddMMAP ./vectorAddMMAP # 使用性能最优设备 ./vectorAddMMAP -device=0 # 显式指定设备编号运行输出示例：设备信息、VIRTUAL ADDRESS MANAGEMENT SUPPORTED = 1、fatbin 加载路径，最终打印Result = PASS（校验失败则打印Result = FAIL并以非零码退出）。
总结
vectorAddMMAP以最简的向量加法为载体，系统展示了 CUDA 虚拟内存管理（VMM）四大步骤：预留 VA（cuMemAddressReserve）→ 创建物理分配（cuMemCreate）→ 映射（cuMemMap）→ 设置访问（cuMemSetAccess），并完整覆盖了释放路径（cuMemUnmap+cuMemAddressFree）。它与vectorAddDrv的对照关系、与 simpleP2P 等 P2P 示例的能力边界，为读者理解"物理属性可定制、虚拟访问仍连续"的现代 CUDA 内存模型提供了可编译、可运行、可逐步断点跟读的实践起点。
【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples
创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
哈希解共线点问题)
