当前位置：首页 > ds >正文

手把手教你用CUDA Graph：将你的LLM推理延迟降低一个数量级

ds 2025/9/5 18:41:38

开篇：一段看似“正常”的“慢”代码

作为一名追求极致性能的开发者，你可能写过类似下面这样的LLM推理服务核心逻辑：它接收请求，将数据拷贝到GPU，执行一系列CUDA核函数，然后将结果拷回。代码清晰，逻辑直接。

C++

// 伪代码: 传统推理循环
for (auto& request : incoming_requests) {// 1. 数据拷贝 D2HcudaMemcpyAsync(d_input, request.data, size, cudaMemcpyHostToDevice, stream);// 2. 执行计算launch_prefill_kernels(d_input, d_kv_cache, ..., stream);for (int i = 0; i < request.gen_length; ++i) {launch_decode_kernel(d_kv_cache, d_output, ..., stream);}// 3. 结果拷回 H2DcudaMemcpyAsync(request.result, d_output, res_size, cudaMemcpyDeviceToHost, stream);cudaStreamSynchronize(stream);
}

这段代码看起来毫无问题，但当你用NVIDIA Nsight等工具剖析时，会发现GPU的计算单元（SM）利用率可能低得惊人，大量时间被耗费在Kernel之间的“空白期”。这就是CPU控制开销的“死亡之谷”。

本文将带你开启一段重构之旅，我们将以上述代码为起点，一步步引入CUDA Graph，让你亲眼见证并理解这场从“命令式”到“声明式”的性能革命是如何发生的。

第一站：诊断性能瓶颈——CPU的“四大原罪”

在重构之前，我们必须精准定位问题。上述循环中的每一行CUDA调用，都隐藏着CPU端的开销：

启动延迟：每一次launch_..._kernel，CPU都要与驱动“沟通”一番，这个固定的“沟通成本”积少成多，对于decode循环中大量的小Kernel来说是致命的。
API开销：cudaMemcpyAsync等调用，同样伴随着CPU与驱动的交互，增加了额外的负担。
调度抖动：这个for循环的执行节奏，完全受控于CPU的操作系统。任何系统中断，都会让GPU的“口粮”供应出现“断档”。
重复劳动：对每个请求，CPU都不厌其烦地重复着几乎完全一样的调用序列，毫无效率可言。

我们的重构目标，就是将CPU从这种繁琐、低效的实时指挥中彻底解放出来。

第二站：首次重构——“录制”我们的第一个计算图

CUDA Graph的核心理念是“一次录制，反复执行”。我们首先将固定的、重复性最高的部分——单步的decode过程，进行“图化”。

C++

// --- 录制阶段 (在服务启动时执行一次) ---
cudaGraph_t decode_graph;
cudaGraphExec_t decode_graph_exec;
cudaStream_t stream = // ... get a stream// 开启“录制模式”
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);// 将你要录制的固定操作序列放入流中
// 注意：此时使用的指针可以是“占位符”，后续可以更新
launch_decode_kernel(placeholder_kv_cache, placeholder_output, ..., stream);
// 你甚至可以包含一些小的数据拷贝等固定操作// 结束录制，生成计算图
cudaStreamEndCapture(stream, &decode_graph);// “实例化”图，使其成为可执行对象，这个过程会进行优化和编译
cudaGraphInstantiate(&decode_graph_exec, decode_graph, NULL, NULL, 0);

通过cudaStreamBeginCapture和cudaStreamEndCapture，我们像使用录像机一样，将launch_decode_kernel这个操作以及它的参数、依赖关系都录制了下来，形成了一个名为decode_graph的蓝图。cudaGraphInstantiate则将这个蓝图编译成了一个高效的、可执行的实例decode_graph_exec。

第三站：见证奇迹——用“一行代码”执行计算

录制完成后，我们原来的decode循环就可以被彻底改写：

C++

// 重构后的推理循环
for (auto& request : incoming_requests) {cudaMemcpyAsync(d_input, request.data, size, cudaMemcpyHostToDevice, stream);launch_prefill_kernels(d_input, d_kv_cache, ..., stream);// 关键改变在这里！for (int i = 0; i < request.gen_length; ++i) {// (在实际执行前，可能需要更新图中的数据指针，详见下一站)update_graph_pointers(decode_graph_exec, d_kv_cache, d_output);// 用一个轻量级的调用，替代原来所有的CPU->GPU指令cudaGraphLaunch(decode_graph_exec, stream);}cudaMemcpyAsync(request.result, d_output, res_size, cudaMemcpyDeviceToHost, stream);cudaStreamSynchronize(stream);
}

看到了吗？原来循环内的launch_decode_kernel调用，变成了一个简单的cudaGraphLaunch。

性能飞跃的秘密：这一行代码，CPU只做了一件事——“告诉GPU，执行你已经知道的那个计划”。GPU接管后，在内部以零开销、背靠背的方式执行所有录制好的操作，彻底绕开了CPU的性能瓶颈。

第四站：应对真实世界的挑战——驾驭动态性

至此，我们已经优化了核心循环。但真实世界的挑战在于动态性。

挑战1：动态输入数据我们的decode_graph在录制时用了占位符指针。在每次cudaGraphLaunch之前，如何让它处理当前请求的真实数据呢？答案是图更新（Graph Update）。

CUDA提供了cudaGraphExecKernelNodeSetParams等接口，允许你在不重新录制图的情况下，高效地更新图中某个节点（例如一个Kernel）的参数。这就实现了“计算流程图固化，输入输出数据动态”的完美结合。

挑战2：动态序列长度 (Prefill阶段) Prefill阶段的计算量与输入Prompt长度相关，是动态的。对此，我们有两种成熟的工程策略：

分段执行（推荐）：承认并接受Prefill阶段的动态性，对这部分不使用CUDA Graph，依然采用传统的Kernel Launch方式。我们只对占总计算步骤95%以上的、固定模式的Decoding阶段进行图化。这是“二八原则”在性能优化中的最佳体现。
装桶策略（Bucketing）：如果输入的序列长度变化范围有限，我们可以为几个典型的长度（如64, 128, 256）预先创建好不同的CUDA Graph。运行时，将请求填充（Padding）到最接近的桶尺寸，然后调用对应的Graph。这是一种经典的“空间换时间”策略。

终点站：集大成者——将一切融入图中

最后，不要把CUDA Graph看作是底层优化的替代品。它是一个更高层次的**“编排者”和“粘合剂”**。在一个高度优化的系统中，你的CUDA Graph节点中，应该包含对cuBLAS、cuDNN等高性能库的调用，以及对FlashAttention这类手工优化的、顶级的自定义Kernel的调用。CUDA Graph负责消除这些强大组件之间的“调用缝隙”，让它们的威力得以100%发挥。