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 APIREADME 列出了该示例使用的全部 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_kernelcuLaunchKernel启动向量加法内核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 doesnt 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_FILEcuModuleLoadData加载二进制模块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输入物理内存需要跨哪些设备条带化驻留stripemappingDevices输入哪些设备需要读写这块内存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.idcuMemCreate(allocationHandle, stripeSize, prop, 0)创建物理分配cuMemMap(*dptr stripeSize * idx, stripeSize, 0, allocationHandle, 0)映射到对应 VA 偏移随后立即cuMemRelease(allocationHandle)——映射建立后句柄不再需要物理内存由映射保持存活。设置跨设备访问权限为每个mappingDevices构造CUmemAccessDesclocation.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 检查的原因。释放实现simpleFreeMultiDeviceMmapmultidevicealloc_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.fatbinfatbin 文件生成后由findFatbinPath在运行时定位加载。可执行文件由vectorAddMMAP.cpp与multidevicealloc_memmap.cpp组成链接CUDA::cuda_driver。该目录通过 cpp/0_Introduction/CMakeLists.txt 中的add_subdirectory(vectorAddMMAP)接入整个仓库的构建体系。典型构建与运行方式在仓库根目录mkdir build cd build cmake .. -DCMAKE_BUILD_TYPERelease make vectorAddMMAP ./vectorAddMMAP # 使用性能最优设备 ./vectorAddMMAP -device0 # 显式指定设备编号运行输出示例设备信息、VIRTUAL ADDRESS MANAGEMENT SUPPORTED 1、fatbin 加载路径最终打印Result PASS校验失败则打印Result FAIL并以非零码退出。总结vectorAddMMAP以最简的向量加法为载体系统展示了 CUDA 虚拟内存管理VMM四大步骤预留 VAcuMemAddressReserve→ 创建物理分配cuMemCreate→ 映射cuMemMap→ 设置访问cuMemSetAccess并完整覆盖了释放路径cuMemUnmapcuMemAddressFree。它与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），仅供参考 cuda-samples vectorAddMMAP 深度解析：用 cuMemMap 虚拟内存管理实现向量加法
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 APIREADME 列出了该示例使用的全部 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_kernelcuLaunchKernel启动向量加法内核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 doesnt 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_FILEcuModuleLoadData加载二进制模块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输入物理内存需要跨哪些设备条带化驻留stripemappingDevices输入哪些设备需要读写这块内存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.idcuMemCreate(allocationHandle, stripeSize, prop, 0)创建物理分配cuMemMap(*dptr stripeSize * idx, stripeSize, 0, allocationHandle, 0)映射到对应 VA 偏移随后立即cuMemRelease(allocationHandle)——映射建立后句柄不再需要物理内存由映射保持存活。设置跨设备访问权限为每个mappingDevices构造CUmemAccessDesclocation.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 检查的原因。释放实现simpleFreeMultiDeviceMmapmultidevicealloc_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.fatbinfatbin 文件生成后由findFatbinPath在运行时定位加载。可执行文件由vectorAddMMAP.cpp与multidevicealloc_memmap.cpp组成链接CUDA::cuda_driver。该目录通过 cpp/0_Introduction/CMakeLists.txt 中的add_subdirectory(vectorAddMMAP)接入整个仓库的构建体系。典型构建与运行方式在仓库根目录mkdir build cd build cmake .. -DCMAKE_BUILD_TYPERelease make vectorAddMMAP ./vectorAddMMAP # 使用性能最优设备 ./vectorAddMMAP -device0 # 显式指定设备编号运行输出示例设备信息、VIRTUAL ADDRESS MANAGEMENT SUPPORTED 1、fatbin 加载路径最终打印Result PASS校验失败则打印Result FAIL并以非零码退出。总结vectorAddMMAP以最简的向量加法为载体系统展示了 CUDA 虚拟内存管理VMM四大步骤预留 VAcuMemAddressReserve→ 创建物理分配cuMemCreate→ 映射cuMemMap→ 设置访问cuMemSetAccess并完整覆盖了释放路径cuMemUnmapcuMemAddressFree。它与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），仅供参考