CUDA Samples 源码解析：graphMemoryFootprint —— 用 CUDA Graph 内存节点复用虚拟地址与物理内存

CUDA Samples 源码解析：graphMemoryFootprint —— 用 CUDA Graph 内存节点复用虚拟地址与物理内存

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples

本篇技术指南以 NVIDIA CUDA Samples 中graphMemoryFootprint示例为核心，系统讲解 CUDA Graphs 的内存分配节点（MemAlloc Node）与内存释放节点（MemFree Node）如何在不同场景下复用虚拟地址（Virtual Address）与物理内存（Physical Memory），并演示如何用cudaDeviceGetGraphMemAttribute/cudaDeviceGraphMemTrim观测与回收图内存池足迹（memory footprint）。读完本文，你将掌握图内存节点的生命周期模型、四种典型内存复用场景的差异，以及如何在单流、多流、未释放分配等情况下精确控制 GPU 显存占用。

示例概览：README 说了什么

cpp/3_CUDA_Features/graphMemoryFootprint/README.md 对该示例的定位非常凝练："This sample demonstrates how graph memory nodes re-use virtual addresses and physical memory."也就是说，它不演示如何用图内存节点做一次性的分配/释放，而是专门研究"复用"——这是图内存节点相对传统cudaMalloc的核心性能优势所在。

README 同时明确了该示例的三大关键概念：CUDA Runtime API、Performance Strategies、CUDA Graphs，以及其支持范围：

维度	支持情况
SM 架构	SM 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、armv7l
运行门槛	驱动支持图内存节点（源码中实际校验驱动 ≥ 11.4.0），且设备支持内存池（cudaDevAttrMemoryPoolsSupported）
前置条件	安装对应平台的 CUDA Toolkit

值得强调的是，README 中给出的18 个 CUDA Runtime API全部可以在 graphMemoryFootprint.cu 中找到真实调用点，本文第 4 节会逐一组装成"API 清单 → 源码调用位置"的对应关系。

构建与运行：CMake 配置与命令行参数

示例通过 CMake 构建，其 CMakeLists.txt 的关键配置如下：

project(graphMemoryFootprint LANGUAGES C CXX CUDA) find_package(CUDAToolkit REQUIRED) set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120) set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Wno-deprecated-gpu-targets") if(ENABLE_CUDA_DEBUG) set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -G") # enable cuda-gdb else() set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -lineinfo") endif() include_directories(../../../Common) add_executable(graphMemoryFootprint graphMemoryFootprint.cu) target_compile_options(graphMemoryFootprint PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>) target_compile_features(graphMemoryFootprint PRIVATE cxx_std_17 cuda_std_17) set_target_properties(graphMemoryFootprint PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

值得注意的实现细节：

• 架构范围：构建目标覆盖 SM 75 到 SM 120（CMAKE_CUDA_ARCHITECTURES），与 README 声明支持的架构范围一致；
• 头文件依赖：通过include_directories(../../../Common)引入仓库 Common 目录，实际使用的是 helper_cuda.h 与 helper_functions.h；
• 编译特性：要求 C++17 / CUDA 17，并开启CUDA_SEPARABLE_COMPILATION；
• 注册方式：该示例通过父级 cpp/3_CUDA_Features/CMakeLists.txt 中的add_subdirectory(graphMemoryFootprint)纳入整个3_CUDA_Features目录的构建体系，因此可随整个仓库统一编译，也可单独cmake构建该子目录。

运行示例时，默认分配大小为64 MiB（源码main中size_t bytes = 64 * 1024 * 1024;），并通过 helper_cuda.h 的findCudaDevice选择设备——支持命令行-device=<id>指定设备号，否则自动选择算力最高（Gflops 最大）的设备。完整的运行方式为：

# 构建（以仓库根目录为例） cmake -S . -B build cmake --build build --target graphMemoryFootprint -j # 运行：自动选择最佳设备 ./build/cpp/3_CUDA_Features/graphMemoryFootprint/graphMemoryFootprint # 运行：指定 GPU 设备 ./build/cpp/3_CUDA_Features/graphMemoryFootprint/graphMemoryFootprint -device=0

运行前置校验：驱动与内存池能力检查

main()在开始任何演示前会做两项硬性校验（graphMemoryFootprint.cu）：

cudaDriverGetVersion(&driverVersion); printf("Driver version is: %d.%d\n", driverVersion / 1000, (driverVersion % 100) / 10); if (driverVersion < 11040) { printf("Waiving execution as driver does not support Graph Memory Nodes\n"); exit(EXIT_WAIVED); } cudaDeviceGetAttribute(&deviceSupportsMemoryPools, cudaDevAttrMemoryPoolsSupported, device); if (!deviceSupportsMemoryPools) { printf("Waiving execution as device does not support Memory Pools\n"); exit(EXIT_WAIVED); }

• 驱动版本门槛：图内存节点依赖驱动提供的 mempool 支持，驱动版本低于11.4.0（11040）时直接以EXIT_WAIVED（其定义见 helper_cuda.h，值为 2）跳过执行，而不是报错；
• 设备能力门槛：通过cudaDeviceGetAttribute(..., cudaDevAttrMemoryPoolsSupported, ...)检查设备是否支持内存池（Memory Pools）。满足条件后，程序依次执行 4 个演示实验，每轮实验结束后调用cleanupMemory()用cudaDeviceGraphMemTrim收回池中空闲内存。

核心 API 全景：README 列出的 18 个运行时 API

下表将 README 的 API 清单与源码中的实际调用位置逐一对应，便于检索与引用：

CUDA Runtime API	作用	源码调用位置
cudaGraphCreate/cudaGraphDestroy	创建/销毁图对象	graphMemoryFootprint.cu
cudaGraphAddMemAllocNode	向图中添加内存分配节点	graphMemoryFootprint.cu
cudaGraphAddMemFreeNode	向图中添加内存释放节点	graphMemoryFootprint.cu
cudaGraphAddKernelNode	向图中添加内核执行节点（用于延长图运行时间）	graphMemoryFootprint.cu
cudaGraphInstantiate/cudaGraphExecDestroy	实例化/销毁可执行图	graphMemoryFootprint.cu
cudaGraphLaunch	在流上启动可执行图	graphMemoryFootprint.cu
cudaStreamCreateWithFlags/cudaStreamDestroy	创建（非阻塞标志）/销毁流	graphMemoryFootprint.cu
cudaStreamSynchronize	等待流中所有工作完成	graphMemoryFootprint.cu
cudaDeviceGetGraphMemAttribute	查询当前图内存池足迹（字节数）	graphMemoryFootprint.cu
cudaDeviceGraphMemTrim	将池中空闲内存归还给设备	graphMemoryFootprint.cu
cudaDeviceGetAttribute	查询设备属性（时钟频率、内存池支持等）	graphMemoryFootprint.cu
cudaGetDeviceProperties/cudaDriverGetVersion	查询设备属性 / 驱动版本	graphMemoryFootprint.cu
cudaFree	释放图外设备内存（配合图内未释放的分配）	graphMemoryFootprint.cu

其中printMemoryFootprint()是贯穿全示例的观测函数（graphMemoryFootprint.cu）：

void printMemoryFootprint(int device) { size_t footprint; checkCudaErrors(cudaDeviceGetGraphMemAttribute(device, (cudaGraphMemAttributeType)0, &footprint)); printf(" FOOTPRINT: %lu bytes\n", footprint); }

它通过属性类型0（即cudaGraphMemAttrUsedMemHighWaterMark之前的cudaGraphMemAttrReservedMemCurrent，表示当前保留的内存字节数）读取设备图内存池的当前足迹。所有实验都以"操作前后分别打印 footprint"的方式量化内存复用效果。

另一个被反复使用的辅助函数prepareAllocParams()（graphMemoryFootprint.cu）统一构造cudaMemAllocNodeParams，其要点是：

allocParams->bytesize = bytes; allocParams->poolProps.allocType = cudaMemAllocationTypePinned; allocParams->poolProps.location.id = device; allocParams->poolProps.location.type = cudaMemLocationTypeDevice;

即：分配bytes字节、采用Pinned（固定）分配类型、物理位置为指定设备（Device）。注意这里allocType使用cudaMemAllocationTypePinned而非cudaMemAllocationTypeMax，意味着图内存节点会从该设备的固定内存池中划拨内存。

实验一：单图内的顺序分配/释放——虚拟地址复用

第一个实验virtualAddressReuseSingleGraph()（graphMemoryFootprint.cu）演示虚拟地址复用。核心逻辑在createVirtAddrReuseGraph()中（graphMemoryFootprint.cu）：

checkCudaErrors(cudaGraphAddMemAllocNode(&allocNodeA, graph, NULL, 0, &allocParams)); d_a = (float *)allocParams.dptr; checkCudaErrors(cudaGraphAddMemFreeNode(&freeNodeA, graph, &allocNodeA, 1, (void *)d_a)); // The dependency between the allocation of d_b and the free of d_a allows d_b // to reuse the same VA. checkCudaErrors(cudaGraphAddMemAllocNode(&allocNodeB, graph, &freeNodeA, 1, &allocParams)); d_b = (float *)allocParams.dptr; if (d_a == d_b) { printf("Check confirms that d_a and d_b share a virtual address.\n"); }

这里的两个关键工程细节：

1. 依赖链驱动复用：d_a的释放节点（freeNodeA）作为d_b分配节点的前驱依赖（&freeNodeA, 1），使得 CUDA 可以在同一个图内、在d_a生命周期结束后把同一段虚拟地址交给d_b；
2. 运行时验证：分配完成后直接比较d_a == d_b指针是否相等。若相等，打印 "d_a and d_b share a virtual address"，即 VA 复用成功。

实验结论正如函数开头打印的说明：单图内顺序的分配与释放，使 CUDA 能够复用虚拟地址。这也解释了图内存节点相对传统内存管理的优势——虚拟地址无需重新映射，分配/释放可在 GPU 工作流内部完成，开销远低于 CPU 发起的cudaMalloc/cudaFree。

实验二：单流上的多图顺序执行——物理内存复用

第二个实验physicalMemoryReuseSingleStream()（graphMemoryFootprint.cu）把视角从"虚拟地址"转向物理内存：8 个彼此独立的可执行图在同一条流上顺序启动。

每个图由createSimpleAllocFreeGraph()构造（graphMemoryFootprint.cu），结构为：MemAlloc 节点 → Kernel 节点（clockBlock）→ MemFree 节点，三者串成依赖链：

checkCudaErrors(cudaGraphAddMemAllocNode(&allocNodeA, graph, NULL, 0, &allocParams)); *dPtr = (float *)allocParams.dptr; int clockRate; checkCudaErrors(cudaDeviceGetAttribute(&clockRate, cudaDevAttrClockRate, device)); clock_t time_clocks = (clock_t)((kernelTime / 1000.0) * clockRate); void *blockDeviceArgs[1] = {(void *)&time_clocks}; ... blockDeviceNodeParams.func = (void *)clockBlock; blockDeviceNodeParams.kernelParams = (void **)blockDeviceArgs; checkCudaErrors(cudaGraphAddKernelNode(&blockDeviceNode, graph, &allocNodeA, 1, &blockDeviceNodeParams)); checkCudaErrors(cudaGraphAddMemFreeNode(&freeNodeA, graph, &blockDeviceNode, 1, (void *)*dPtr));

这里的内核节点clockBlock（graphMemoryFootprint.cu）是一个"空转"内核——它不做实际计算，只通过clock()忙等至少clock_count个时钟周期。其作用（源码注释明确说明）是延长每个图的运行时间，从而保证多个图在同一条流中不会因为执行过快而重叠生命周期，确保"顺序执行、生命周期不重叠"的测试前提成立。

实验通过 4 次 footprint 打印，定量展示了物理内存复用的完整过程：

1. 创建图执行体不占物理内存：循环创建 8 个graphExec后打印，footprint 不增长（Creating the graph execs does not reserve any physical memory.）；
2. 首次启动才保留内存：cudaGraphLaunch(graphExecs[0], stream)后 footprint 增长——"The first graph launched reserves the memory it needs."；
3. 同图重复启动复用内存：再次启动graphExecs[0]，footprint 不再增长；
4. 不同图顺序启动同样复用：依次启动graphExecs[1..7]，footprint 始终保持不变——"Subsequent launches of other graphs in the same stream also reuse the physical memory."

实验末尾还做了一项正确性验证：两两比较 8 个图的dPtrs[i]，确认所有图使用不同的虚拟地址（virtualAddrDiffer保持为 true），以证明"物理内存被复用、但虚拟地址各不相同"，即复用发生在物理页层面而非 VA 层面。

实验三：并发流上的多图并行——内存足迹必然增长

第三个实验simultaneousStreams()（graphMemoryFootprint.cu）探讨相反的场景：能够并发运行的图必须各自持有独立的物理内存。

for (int i = 0; i < NUM_GRAPHS; i++) { createSimpleAllocFreeGraph(&graphExecs[i], &dPtrs[i], bytes, device); checkCudaErrors(cudaStreamCreateWithFlags(&streams[i], cudaStreamNonBlocking)); } ... for (int i = 1; i < NUM_GRAPHS; i++) { checkCudaErrors(cudaGraphLaunch(graphExecs[i], streams[i])); printf("%02d: ", i); printMemoryFootprint(device); }

8 个图分别被发射到 8 条独立的非阻塞流上，彼此可能同时执行。由于各图的生命周期重叠，CUDA 无法把同一块物理内存同时租给两个并发图，因此每次启动都会观察到 footprint 单调增长。

源码注释同时指出了 CUDA 的一个精细行为：启动新图时，CUDA 可能复用"已经执行完毕"的图的物理内存——即便新图位于不同流。这正是示例在内核节点中加入clockBlock空转内核的原因（注释原文："a kernel node is added to the graphs to increase runtime"）：如果各图运行太快，后面的图可能复用前面已结束图的内存，从而"测不出"真实的并发内存占用。延长运行时间确保所有图真正处于并发执行状态，让内存复用无法发生，从而如实展示并发场景下的足迹增长。

实战启示：如果你的工作负载需要多个图并发执行，内存规划必须按"并发图数量 × 单图峰值分配"估算显存；反之，若图只在单流中顺序执行，则物理内存可以按单图峰值近似估算。

实验四：未释放的分配与内存裁剪（Trim）

第四个实验unfreedAllocations()（graphMemoryFootprint.cu）研究错误/延迟释放对内存足迹的影响。它使用createSimpleAllocNoFreeGraph()（graphMemoryFootprint.cu）构造只有 MemAlloc 节点、没有 MemFree 节点的图：

checkCudaErrors(cudaGraphAddMemAllocNode(&allocNodeA, graph, NULL, 0, &allocParams)); *dPtr = (float *)allocParams.dptr; checkCudaErrors(cudaGraphInstantiate(graphExec, graph, NULL, NULL, 0));

实验结果分三个阶段，每一步都用 footprint 打印佐证：

1. 同一流顺序启动也不复用：8 个未释放分配的图即使在同一流上顺序启动，footprint 也持续增长。原因正如注释所说："Since the allocation is not freed, CUDA keeps the memory valid for use."——内存必须保持有效以备后续使用，无法回收；
2. Trim 无法回收有效内存：调用cudaDeviceGraphMemTrim(device)后 footprint 不变。因为 trim 只能把池中空闲的内存归还设备，而"已被分配且未释放"的内存仍处于使用中；
3. 释放后 Trim 才生效：先用cudaFree(dPtrs[i])释放全部指针（注意：此时图执行体graphExec尚未销毁），footprint 仍不下降（内存回到池中但未归还设备）；随后再次调用cudaDeviceGraphMemTrim(device)，footprint 显著下降——"Since the allocations are now freed, trimming does reduce the footprint even when the graph execs are not yet destroyed."

cleanupMemory()（graphMemoryFootprint.cu）在每个实验结束后执行同样的 trim 操作，确保进入下一实验前设备内存已被清理，避免实验间相互干扰：

void cleanupMemory(int device) { checkCudaErrors(cudaDeviceGraphMemTrim(device)); printf("\nCleaning up example by trimming device memory.\n"); printMemoryFootprint(device); }

内存管理三阶段模型（从本实验可以提炼出图内存池的完整生命周期）：cudaGraphLaunch时按需从设备预留物理内存（footprint 上升）→ 图内 MemFree 节点释放后内存返回池中（footprint 不变）→cudaDeviceGraphMemTrim将池中空闲内存真正归还设备（footprint 下降）。理解这一模型，才能在长生命周期应用中主动控制显存占用，避免"峰值显存虚高"。

四个实验的对照总结与性能策略

实验	场景	生命周期关系	内存复用	footprint 表现
实验一	单图内顺序 alloc/free	串行、不重叠	虚拟地址复用（d_a == d_b）	不增长
实验二	多图、单流顺序启动	串行、不重叠	物理内存复用（VA 各异）	首次启动后不再增长
实验三	多图、多流并发启动	并行、重叠	无法复用（各图独立物理内存）	每个图启动都增长
实验四	多图、无 MemFree 节点	分配后长期持有	无法复用，trim 也无效	持续增长，释放+trim 后才下降

综合来看，graphMemoryFootprint为使用 CUDA Graph 内存节点的开发者给出了三条可落地的性能策略：

1. 尽量让分配的生命周期在单图内闭合（配 MemFree 节点），以便 CUDA 在同一图内复用虚拟地址，省去 VA 重映射开销；
2. 把生命周期不重叠的图调度到同一条流上顺序执行，让它们共享同一份物理内存，把峰值显存压到"单图峰值"量级；
3. 用cudaDeviceGetGraphMemAttribute定期观测足迹、用cudaDeviceGraphMemTrim主动回收池中空闲内存，尤其警惕"只分配不释放"的图长期驻留导致显存膨胀。

如需进一步对比，仓库中的姊妹示例 graphMemoryNodes（图内存节点的分配/释放演示，同时覆盖 Graph API 与 Stream Capture 两种构建方式）与 graphConditionalNodes、graphMemoryFootprint 同属3_CUDA_Features目录，可配合阅读，完整理解 CUDA Graphs 的节点能力矩阵。

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples