cuda-samples 之 cudaCompressibleMemory基于 cuMemMap 与 Driver API 的可压缩内存分配实战解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文围绕 NVIDIA CUDA Samples 仓库中的cudaCompressibleMemory示例展开深入剖析如何通过 CUDA Driver API 的虚拟内存管理接口cuMemMap系列 API申请可压缩内存Compressible Memory并在其上执行 SAXPY 向量运算以验证读写与性能特性。读完本文你将掌握CUmemAllocationProp配置、cuMemCreate/cuMemMap/cuMemSetAccess的完整分配流程、设备压缩能力检测、不同 SM 架构如 SM 8.0/8.6下的初始化限制以及该示例的构建与运行方式。示例概览与仓库定位cudaCompressibleMemory位于 cpp/3_CUDA_Features/cudaCompressibleMemory 目录是 CUDA Features 章节下的一个 Driver API 示例用于演示使用cuMemMapAPI 进行可压缩内存分配。其核心文件构成如下README.md示例说明与 API 清单。compMalloc.cpp可压缩/普通内存的分配与释放封装allocateCompressible/freeCompressible。compMalloc.h上述两个接口的声明。saxpy.cu入口程序负责设备能力检测、内存分配与 SAXPY 内核启动。CMakeLists.txt该示例的构建脚本。该示例通过add_subdirectory(cudaCompressibleMemory)被纳入 cpp/3_CUDA_Features/CMakeLists.txt 的构建体系。在运行时程序会分别用可压缩与不可压缩内存各执行一次 SAXPY 并报告耗时与带宽方便对比两种内存类型的实际表现。关键概念什么是可压缩内存可压缩内存是 CUDA 提供的一种内存分配类型当数据本身具备较强的可压缩性例如大量重复数值、稀疏结构或零值较多的缓冲区时硬件可在显存中对其进行压缩存储从而提升有效内存容量并可能降低显存带宽压力进而改善整体性能。需要强调的是并非所有数据都适合压缩只有对压缩友好的数据才可能获得收益可压缩内存的分配与使用必须依赖虚拟内存管理VMM能力这也是本示例选用cuMemMap系列 Driver API 的根本原因压缩是透明的内核代码无需感知压缩过程按普通全局内存读写即可如本示例的 SAXPY 内核。支持的平台与架构根据 README.md 的说明该示例支持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、QNXCPU 架构x86_64、aarch64。不过是否真正能使用可压缩内存最终取决于运行时设备的实际能力程序会在启动阶段通过设备属性查询来判定详见下文设备能力检测一节。从 CMakeLists.txt 的默认编译目标架构75 80 86 87 89 90 100 110 120可以看出当前仓库构建脚本面向的是较新的计算能力版本。涉及的核心 CUDA APIDriver APIAPI作用cuCtxGetDevice获取当前上下文关联的设备用于配置内存位置cuDeviceGetAttribute查询设备属性VMM 支持、压缩支持、计算能力cuMemGetAllocationGranularity查询内存分配粒度用于将请求大小向上对齐cuMemAddressReserve预留一段虚拟地址范围cuMemCreate根据分配属性创建物理内存分配可压缩cuMemGetAllocationPropertiesFromHandle校验实际分配到的内存属性确认压缩类型生效cuMemMap将物理分配映射到预留的虚拟地址cuMemRelease释放物理分配句柄引用cuMemSetAccess设置映射内存的访问权限cuMemUnmap/cuMemAddressFree解除映射并释放虚拟地址范围Runtime APIcudaMemcpy压缩缓冲区的初始化、cudaEventCreate/cudaEventRecord/cudaEventSynchronize/cudaEventElapsedTime性能计时、cudaOccupancyMaxPotentialBlockSize自动计算内核的 block 大小与 grid 大小。内存分配属性配置setPropsetPropcompMalloc.cpp负责填充CUmemAllocationProp结构体是可压缩分配的核心配置cudaError_t setProp(CUmemAllocationProp *prop, bool UseCompressibleMemory) { CUdevice currentDevice; if (cuCtxGetDevice(currentDevice) ! CUDA_SUCCESS) return cudaErrorMemoryAllocation; memset(prop, 0, sizeof(CUmemAllocationProp)); prop-type CU_MEM_ALLOCATION_TYPE_PINNED; prop-location.type CU_MEM_LOCATION_TYPE_DEVICE; prop-location.id currentDevice; if (UseCompressibleMemory) prop-allocFlags.compressionType CU_MEM_ALLOCATION_COMP_GENERIC; return cudaSuccess; }要点说明prop-type CU_MEM_ALLOCATION_TYPE_PINNED分配类型为固定设备端内存prop-location.type CU_MEM_LOCATION_TYPE_DEVICE与prop-location.id currentDevice指定分配所在的设备prop-allocFlags.compressionType CU_MEM_ALLOCATION_COMP_GENERIC请求通用压缩类型——这是开启可压缩内存的关键开关当UseCompressibleMemory为false时该字段保持为 0即普通不可压缩分配。可压缩内存分配流程allocateCompressibleallocateCompressiblecompMalloc.cpp完整演示了分配粒度对齐 → 预留虚拟地址 → 创建物理分配 → 校验压缩类型 → 建立映射 → 设置访问权限的 VMM 标准流程cudaError_t allocateCompressible(void **adr, size_t size, bool UseCompressibleMemory) { CUmemAllocationProp prop {}; cudaError_t err setProp(prop, UseCompressibleMemory); if (err ! cudaSuccess) return err; size_t granularity 0; if (cuMemGetAllocationGranularity(granularity, prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM) ! CUDA_SUCCESS) return cudaErrorMemoryAllocation; size ((size - 1) / granularity 1) * granularity; CUdeviceptr dptr; if (cuMemAddressReserve(dptr, size, 0, 0, 0) ! CUDA_SUCCESS) return cudaErrorMemoryAllocation; CUmemGenericAllocationHandle allocationHandle; if (cuMemCreate(allocationHandle, size, prop, 0) ! CUDA_SUCCESS) return cudaErrorMemoryAllocation; // Check if cuMemCreate was able to allocate compressible memory. if (UseCompressibleMemory) { CUmemAllocationProp allocationProp {}; cuMemGetAllocationPropertiesFromHandle(allocationProp, allocationHandle); if (allocationProp.allocFlags.compressionType ! CU_MEM_ALLOCATION_COMP_GENERIC) { printf(Could not allocate compressible memory... so waiving execution\n); exit(EXIT_WAIVED); } } if (cuMemMap(dptr, size, 0, allocationHandle, 0) ! CUDA_SUCCESS) return cudaErrorMemoryAllocation; if (cuMemRelease(allocationHandle) ! CUDA_SUCCESS) return cudaErrorMemoryAllocation; CUmemAccessDesc accessDescriptor; accessDescriptor.location.id prop.location.id; accessDescriptor.location.type prop.location.type; accessDescriptor.flags CU_MEM_ACCESS_FLAGS_PROT_READWRITE; if (cuMemSetAccess(dptr, size, accessDescriptor, 1) ! CUDA_SUCCESS) return cudaErrorMemoryAllocation; *adr (void *)dptr; return cudaSuccess; }各步骤的工程细节粒度对齐通过cuMemGetAllocationGranularity(..., CU_MEM_ALLOC_GRANULARITY_MINIMUM)查询最小分配粒度再把请求大小向上取整对齐size ((size - 1) / granularity 1) * granularity。VMM 分配的尺寸必须是粒度的整数倍否则后续调用会失败。预留虚拟地址cuMemAddressReserve在 GPU 地址空间中预留一段连续虚拟地址本例不指定固定起始地址与对齐约束。创建物理分配cuMemCreate依据prop创建物理内存CUmemGenericAllocationHandle是对该物理分配的句柄。校验压缩是否生效请求压缩后用cuMemGetAllocationPropertiesFromHandle回读句柄的实际属性确认allocFlags.compressionType CU_MEM_ALLOCATION_COMP_GENERIC。若设备未能满足压缩请求例如驱动不支持程序会打印提示并调用exit(EXIT_WAIVED)该退出码 2 定义于 Common/helper_cuda.h主动豁免执行避免在非预期环境下运行。建立映射cuMemMap将物理分配映射到预留的虚拟地址之后cuMemRelease释放句柄引用映射仍有效物理内存生命周期由映射管理。设置访问权限cuMemSetAccess通过CUmemAccessDesc设备位置 CU_MEM_ACCESS_FLAGS_PROT_READWRITE读写权限使映射对设备可访问最终把虚拟地址返回给调用方。内存释放流程freeCompressiblefreeCompressiblecompMalloc.cpp是分配流程的逆操作同样先做粒度对齐然后依次调用cuMemUnmap解除映射、cuMemAddressFree释放虚拟地址范围若传入空指针则直接返回成功。if (ptr NULL) return cudaSuccess; if (cuMemUnmap((CUdeviceptr)ptr, size) ! CUDA_SUCCESS || cuMemAddressFree((CUdeviceptr)ptr, size) ! CUDA_SUCCESS) return cudaErrorInvalidValue; return cudaSuccess;这里体现了 VMM 的生命周期管理原则虚拟地址与物理分配是分离的资源映射时由cuMemMap绑定释放时先解绑cuMemUnmap再回收地址空间cuMemAddressFree物理内存则由分配句柄机制自动回收。设备能力检测与架构差异处理saxpy.cu的main函数saxpy.cu在分配内存前做了两级能力检测虚拟内存管理支持查询CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED若设备不支持 VMM 则打印提示并以EXIT_WAIVED退出通用压缩支持查询CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED若为 0 则同样豁免执行通过后打印Generic memory compression support is available。随后读取CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR/MINOR获取计算能力用于处理一个重要的架构差异saxpy.cubool compressibleZbuf 0; if ((major 8 minor 0) || (major 8 minor 6)) { // On SM 8.0 and 8.6 GPUs compressible buffer can only be initialized // through cudaMemcpy. printf(allocating non-compressible Z buffer\n); checkCudaErrors(allocateCompressible((void **)z, size, false)); compressibleZbuf 0; } else { checkCudaErrors(allocateCompressible((void **)z, size, true)); compressibleZbuf 1; }关键结论在 SM 8.0如 A100与 SM 8.6如 RTX 30 系列架构上压缩缓冲区只能通过cudaMemcpy初始化不能由内核直接写入。因此程序在这些架构上把输出缓冲区z分配为不可压缩内存并通过launchSaxpy内部的compressibleZbuf标志决定初始化路径compressibleZbuf 0在主机端构造h_x/h_y数组并用cudaMemcpy(HostToDevice)拷贝初始化x、ysaxpy.cu其他架构直接调用内核init在设备端初始化saxpy.cu。SAXPY 内核与性能度量SAXPY 内核saxpy.cu使用float4向量化访问每次处理 4 个浮点分量计算z a * x y并通过gridDim.x * blockDim.x步长实现网格跨越grid-stride循环适配任意数据规模__global__ void saxpy(const float a, const float4 *x, const float4 *y, float4 *z, const size_t n) { for (size_t i blockIdx.x * blockDim.x threadIdx.x; i n; i gridDim.x * blockDim.x) { const float4 x4 x[i]; const float4 y4 y[i]; z[i] make_float4(a * x4.x y4.x, a * x4.y y4.y, a * x4.z y4.z, a * x4.w y4.w); } }性能度量saxpy.cu使用事件计时通过cudaOccupancyMaxPotentialBlockSize自动计算minGridSize与blockSize兼顾占用率与可移植性用cudaEventCreate/cudaEventRecord/cudaEventSynchronize/cudaEventElapsedTime测量内核耗时按(size * 3) / ms / 1e9估算有效带宽TB/s其中size * 3对应读x、读y、写z三个方向的流量默认数据规模为n 10485760saxpy.cu即每缓冲区约 40 MBfloat4。程序随后分别以可压缩true与不可压缩false两种方式分配x、y、z并各运行一次 SAXPY打印两轮计时结果供用户对比压缩内存的带宽表现。文件末尾saxpy.cu还附有一句重要提示CUDA Samples 并非为性能测量设计GPU Boost 开启时结果可能波动。命令行参数程序支持-devicenn 0指定运行设备if (checkCmdLineFlag(argc, (const char **)argv, help) || checkCmdLineFlag(argc, (const char **)argv, ?)) { printf(Usage -devicen (n 0 for deviceID)\n); exit(EXIT_SUCCESS); } findCudaDevice(argc, (const char **)argv);findCudaDevice是 Common/helper_cuda.h 提供的公共工具函数负责解析-device并完成设备选择未指定时自动选择可用设备。运行示例./cudaCompressibleMemory -device0构建方式该示例可通过仓库根目录的 CMake 体系构建其独立 CMakeLists.txt 要点如下project(cudaCompressibleMemory LANGUAGES C CXX CUDA)声明语言支持find_package(CUDAToolkit REQUIRED)查找 CUDA Toolkit目标由compMalloc.cpp与saxpy.cu组成add_executable(cudaCompressibleMemory compMalloc.cpp saxpy.cu)开启CUDA_SEPARABLE_COMPILATION启用cxx_std_17/cuda_std_17标准通过target_link_libraries(... CUDA::cuda_driver)链接 Driver API 库本示例是纯 Driver API 用法头文件路径包含仓库 Common 目录用于helper_cuda.h等公共工具并引入 cmake/InstallSamples.cmake 的安装配置。在仓库根目录执行标准 CMake 流程即可构建mkdir build cd build cmake .. make cudaCompressibleMemory运行前提与豁免机制小结综合上文成功运行并真正使用可压缩内存需要满足以下前提否则程序会以EXIT_WAIVED值 2主动退出而非报错设备支持虚拟内存管理CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED设备支持通用压缩CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTEDcuMemCreate实际返回的分配属性满足CU_MEM_ALLOCATION_COMP_GENERIC。这种探测—回退—豁免的设计是 Driver API 特性类示例的通用模式既能覆盖多种硬件配置也保证了程序在旧设备上的可移植性。对于开发者而言可压缩内存适合缓存类、稀疏矩阵、重复数据等压缩友好场景在引入自己的应用中时建议参考本示例的流程先做能力检测再通过cuMemGetAllocationPropertiesFromHandle确认压缩真正生效并注意 SM 8.0/8.6 上压缩缓冲区只能经cudaMemcpy初始化的约束。总结cudaCompressibleMemory是一个麻雀虽小五脏俱全的 Driver API 示例从CUmemAllocationProp的属性配置到粒度对齐、虚拟地址预留、物理分配、映射、权限设置、属性回读校验的完整 VMM 流程再到设备能力检测与架构差异分支处理以及基于事件计时的带宽评估覆盖了将可压缩内存落地到真实内核的全部关键环节。它既可以作为学习cuMemMap系列 API 的入门范本也可以作为评估特定设备压缩内存能力的可复用代码骨架——allocateCompressible与freeCompressible这两个封装函数可直接抽取到自有项目中。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考