直接调用kernel()因每次触发完整主机端调度(参数序列化、入队、同步检查、PCIe下发),小核或高频调用时开销可达10%–20%;CUDA Graph通过捕获单流中操作序列构建静态图,实例化后一次启动即可复用,显著降低launch开销。

为什么直接用 kernel>>() 会拖慢重复计算
每次调用三尖括号语法,CUDA 都要走完整套主机端调度路径:参数序列化、命令入队、流同步检查、PCIe 指令下发。对固定结构的计算流程(比如推理中每帧都跑相同的 layer 序列),这部分开销可能占到总执行时间的 10%–20%,尤其在小 kernel 或高频率调用时更明显。
典型症状包括:nvprof 显示大量 cudaLaunchKernel 时间占比高、GPU 利用率曲线出现密集毛刺、cudaEventRecord 测得两次 launch 间隔远大于 kernel 实际运行时间。
用 cudaGraphCreate 构建静态图的最小可行步骤
核心是把“分配内存 → 拷贝输入 → 启动 kernel → 拷贝输出”这一整条链,打包成一个图对象,后续只调用一次启动指令。
- 先用
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal)开启捕获,所有在此 stream 中发生的操作(cudaMemcpyAsync、kernel>>、cudaMemPrefetchAsync)都会被记录 - 执行一遍完整流程(注意:此时不真正执行,只是录制依赖关系)
- 调用
cudaStreamEndCapture(&stream, &graph)得到图对象 - 用
cudaGraphInstantiate(&instance, graph, nullptr, nullptr, 0)创建可执行实例(这时才做参数绑定和资源预分配)
关键点:捕获必须在**单个 stream** 中完成;不能在捕获期间调用 cudaDeviceSynchronize() 或跨 stream 同步;所有 kernel 启动必须使用该 stream。
立即学习“C++免费学习笔记(深入)”;
cudaGraphUpdate 能改什么、不能改什么
CUDA 10.0+ 支持有限更新,但不是“动态图”。它只允许修改已存在节点的参数,不能增删节点或改变拓扑。
- 能安全更新的:
cudaKernelNodeParams::func(换另一个已注册的 kernel)、cudaKernelNodeParams::gridDim和blockDim、cudaKernelNodeParams::kernelParams(即传给 kernel 的指针参数) - 不能更新的:
cudaMemcpyNodeParams::dst或src地址类型(比如从 host 换成 device)、节点执行顺序、新增一个 memcpy 节点 - 更新前必须确保新参数地址已映射到图上下文(例如用
cudaGraphExecMemcpyNodeSetParams而非直接改结构体字段)
常见误用:试图用 cudaGraphUpdate 改变图结构来适配不同 batch size——这会导致 cudaErrorInvalidValue,应改用多个预构建图 + 条件选择。
CUDA Graph 在 C++ 中如何与 RAII 结合
图对象(cudaGraph_t)、实例(cudaGraphExec_t)都是裸指针,C++ 项目里容易漏掉销毁导致显存泄漏或句柄耗尽。
建议封装轻量 wrapper:
struct CudaGraph {
cudaGraph_t graph = nullptr;
cudaGraphExec_t exec = nullptr;
~CudaGraph() {
if (exec) cudaGraphExecDestroy(exec);
if (graph) cudaGraphDestroy(graph);
}
};注意:cudaGraphDestroy 不会自动销毁关联的 exec,必须显式调用 cudaGraphExecDestroy;反之,销毁 exec 不影响原图。
真正容易被忽略的是:图实例(exec)在多线程环境下不是线程安全的。如果多个线程共用同一个 exec 并发调用 cudaGraphLaunch,需加锁或为每个线程分配独立实例。


















