llama.cpp 源码完全拆解（二十三）：graph_compute 的图路径——键、uid、分块捕获与 cudaGraphExecUpdate 的逐节点比对

一次 decode 步的 CPU 侧开销

解码阶段每一步只新增一个 token，GPU 上跑的是小矩阵乘向量、归一化、注意力这些短 kernel，单个 kernel 在 GPU 上的执行时间常常只有几微秒，而 CPU 侧要依次完成：Python/C++ 层的参数准备（这里 llama.cpp 没有 Python，但模板展开、图重排、张量视图仍有代价）、CUDA runtime 的 cudaLaunchKernel 参数打包、驱动层的工作队列提交。一个 70 层左右的模型，每层十余个算子，一步就是上千次 launch。即便每次 launch 只花 35 微秒，累计也有 35 毫秒，足以让 GPU 在每个 kernel 之间不断空转。

CUDA Graph 的核心是把「一串 kernel 及其依赖关系」一次性录制成一张有向无环图，之后每次只需要一次 cudaGraphLaunch 把整段工作提交给 GPU。CUDA Graph 最早随 CUDA 10（2018 年发布）引入，核心 API 包括 cudaStreamBeginCapture/cudaStreamEndCapture（流捕获）与 cudaGraphInstantiate/cudaGraphLaunch（实例化与回放）；PyTorch 在 1.10 版本（2021 年）通过 stream capture 暴露了这一能力，其官方博客在 2021 年给出了一个直观的时间线对比。

图源：[PyTorch 官方博客 Accelerating PyTorch with CUDA Graphs](https://pytorch.org/blog/accelerating-pytorch-with-cuda-graphs/)（Figure 1）。空窗的长度就是 CPU 提交开销，图化之后这些空窗被压缩成一次提交。

难点：llama.cpp 的执行形态和 PyTorch 不一样

PyTorch 的典型图化用法是：准备好固定形状的静态输入张量，用示例输入先 eager 跑几轮热身，然后 with torch.cuda.graph(g): 在捕获流里跑一次，之后每次把新数据 copy_ 进静态输入再 replay()。形状变了就重新捕获一个新图，或者用 make_graphed_callables 为每个可捕获模块单独建图。这种模式依赖框架在 Python 层做「按形状分桶」。

llama.cpp 的路径则完全不同，它没有「每档 batch 一个静态图」的显式管理，而是让 GPU 后端在每次 graph_compute 时自己去判断「这张图和上次那张是不是同一张」：

• 计算图常常每步都重建，但重建是「确定性的」。解码步走 llama_context::process_ubatch：当上一张图仍活跃、且 llm_graph_result::can_reuse(gparams) 返回真时，直接复用旧图，只通过 set_inputs(ubatch) 刷新输入张量的内容；否则先 res->reset()、ggml_backend_sched_reset(sched)，再 model.build_graph(gparams) 重建。reset() 的关键在于它并不释放元数据内存，而是把 buf_compute_meta 原样传给 ggml_init，再 ggml_new_graph_custom 开一张新图——新张量依旧从同一段 buffer 的开头依次分配，因此张量结构体指针会重复出现。这正是「首节点地址可以当图键」的前提；真正会变的是 data 指针（调度器分配的 buffer 可能换位置）与 ne/nb（batch 大小不同），这也正是逐节点快照要记录的内容。代码里对复用路径有一句很精确的注释：为了正确复用，图的完整拓扑必须由这些图参数唯一决定——同样的约束也间接约束了后端侧的图身份比对。
• 图可能被调度器切成多段。ggml_backend_sched_split_graph 把大图按后端切分成若干 split（比如中间夹着一段 CPU 上的 MoE 专家），每个 split 用 ggml_graph_view 变成子图再交给对应后端。于是同一份计算在不同后端上会变成不同首节点的子图，各自需要独立的捕获对象。
• 同一个 CUDA 上下文里有多条图在轮流出现。生成阶段（tg，每序列一步一个 token）与提示处理阶段（pp，一次处理若干 token）的图结构不同；can_reuse 只看图参数（ubatch 形态、图类型、记忆上下文等）是否一致，序列数、输出数一变就要重建。混合 MoE 的 --n-cpu-moe 部署下，每层还会被 CPU 段打断成多个 GPU split。所有这些形态都要在同一份 cuda_graphs 台账里区分清楚、互不串用。

所以 llama.cpp 必须自己维护一套「图身份」的判定，而且判定粒度不能太粗（粗了会回放错误的图），也不能太细（细了就每步都重新捕获，图化白做）。下面这张 Nsight 时间线能说明「太细」的代价：当 CPU 被 launch 逻辑占满、GPU 在 kernel 之间大量空转时，省下来的空隙本来就不多，若每步都重新捕获，捕获自身的开销很容易把收益吃回去。

图源：[PyTorch 官方博客 Accelerating PyTorch with CUDA Graphs](https://pytorch.org/blog/accelerating-pytorch-with-cuda-graphs/)（Figure 3，Mask R-CNN 的 Nsight 时间线）。同样的形态在解码阶段更极端：每一步的 kernel 更短、更密，launch 开销占比更高。

PyTorch 官方博客在 2021 年给出的数据可以当作量级参考：在 Mask R-CNN 上，只把 resnet 主干图化，被图化的那一段从 31 ms 降到 6 ms（约 5 倍），整模型端到端约 1.7 倍；BERT 在 4096 卡配置下整体约 1.12 倍；推荐模型 DLRM 的收益则在小 batch 下更明显，因为小 batch 时单个 kernel 更短、CPU 提交在总时间里的占比更高。这些是训练负载的数据，但机制完全一致：批量越小、kernel 越碎，launch 开销占比越高，图化的收益越大。解码阶段的每步前向正是「批量小、kernel 碎」的极端情形。

图源：[PyTorch 官方博客 Accelerating PyTorch with CUDA Graphs](https://pytorch.org/blog/accelerating-pytorch-with-cuda-graphs/)（Figure 6，DLRM 的基准结果）。横轴为 batch size，可以看到收益随批量减小而放大——这正是解码阶段最需要图化的原因。

图化的三条硬约束

在进入代码之前，先把 CUDA Graph 对图内节点提出的要求摆出来，后面每一种特殊处理都能对号入座：

1. 参数地址静态。图里每个 kernel 读写的 buffer 地址在捕获时确定，回放时不会变。所以本次计算必须把结果写进「上次捕获时用的同一批地址」。
2. 形状静态。一个 kernel 的 block/grid 维度、向量化宽度等都是从张量形状推出来的，形状变了意味着必须换一个 kernel 配置，旧图里记下的配置无法使用。
3. 图内不允许隐含同步。捕获期间不允许 CPU 与 GPU 之间同步、不允许访问主机端可页交换内存、不允许多个流之间的无序交互（除非用事件显式建模）。llama.cpp 的 MoE mul_mat_id 在某些硬件上要走「先算 ids、回读、再排序」的 fallback 路径，其中包含 cudaStreamSynchronize；一旦图里出现这种节点，整张图就不能捕获。

数据结构：ggml_cuda_graph 里到底存了什么

这个结构定义在 ggml/src/ggml-cuda/common.cuh 中，并且只在编译开启了 GGML_CUDA_USE_GRAPHS（HIP 为 GGML_HIP_GRAPHS）时才有成员：

1
struct ggml_cuda_graph {
2
#ifdef USE_CUDA_GRAPH
3
    ~ggml_cuda_graph() {
4
        for (cudaGraphExec_t instance : instances) {
5
            CUDA_CHECK(cudaGraphExecDestroy(instance));
6
        }
7
    }
8
    // the ggml graph is captured as a series of CUDA graphs,
9
    // so the GPU can run one chunk while the next is enqueued
10
    std::vector<cudaGraphExec_t> instances;
11
    size_t num_nodes = 0;
12
    std::vector<cudaGraphNode_t> nodes;
13
    bool disable_due_to_gpu_arch = false;
14
    bool warmup_complete = false;
15
    uint64_t uid = 0;
16
    int64_t last_used_time = 0;
17
    struct node_properties {
18
        ggml_tensor node;
19
        void *   node_src_data_ptrs[GGML_MAX_SRC];
20
        int64_t  node_src_ne[GGML_MAX_SRC][GGML_MAX_DIMS];
21
        size_t   node_src_nb[GGML_MAX_SRC][GGML_MAX_DIMS];
22
    };
23
    std::vector<node_properties> node_props;
24

25
    bool is_enabled() const {
26
        static const bool disable_cuda_graphs_due_to_env =
27
            (getenv("GGML_CUDA_DISABLE_GRAPHS") != nullptr);
28
        return !(disable_due_to_gpu_arch || disable_cuda_graphs_due_to_env);
29
    }
30
#endif
31
};

逐字段读：

• instances：这是「一张 ggml 图 → 多张 CUDA 图」映射的产物。块大小由 GGML_CUDA_GRAPH_CHUNK_SIZE 决定，默认 512 个 ggml 节点一块；为什么分块、分块的边界为什么不能随便切，见后面的小节。
• node_props：同一性比对的台账。注意它不是只存「指针」或只存「形状」，而是把三个维度都存了：整个 ggml_tensor 结构体（包含 op、op_params、flags 等）、每个源指针 src[j]->data、每个源的形状 ne 与步长 nb。这一组合基本覆盖了「同一张图」的所有可观察差异。
• warmup_complete：热身标志，参与「要不要开始捕获」的判定。下一小节细讲。
• uid：上一次被认定「与这张图相同」的 ggml 图 uid（由 ggml_graph_next_uid() 全局递增分配）。这是判断「要不要做逐节点比对」的快捷通道。
• last_used_time 与外部淘汰扫描：图对象按需创建，也按需回收——最近 10 秒没被用到的图会被清掉，避免大量不同的图把显存里的捕获资源一直占着。

ggml_cuda_graph 对象本身由上下文里的 std::unordered_map<const void *, std::unique_ptr<ggml_cuda_graph>> cuda_graphs 持有，键是「首节点指针」。这个选择很讲究：跨后端切分后，每个 split 都是 ggml_graph_view，它的 nodes[0] 指向原大图里一段连续节点的第一个；在重建确定性的前提下，这个指针会重复出现。为什么不直接用 split 的序号（或 split.graph 的地址）做键？因为调度器重新切图后 split 的数量与次序都可能变——比如某些层被移出 GPU、或者 MoE 段的边界挪了位置——同一个序号会指到完全不同的计算段，旧图就会被错配。首节点是「计算段内容」的一部分，跟着内容走，比跟着容器走更可靠。

第一道关卡：graph_key 到底指向谁

1
static const void * ggml_cuda_graph_get_key(ggml_cgraph * cgraph) {
2
    return cgraph->nodes[0];
3
}

只有一行，却决定了这条路径在混合 CPU/GPU 部署下是否成立。代码里注释写得很直接：「Map from first_node_ptr to cuda_graph - allows multiple graphs per context when the computation is split across CPU/GPU (e.g., with --n-cpu-moe)」。

设想一个开着 --n-cpu-moe 的模型，专家段被放在 CPU 上：第 1 层是 [GPU: attention + FFN]；第 2 层是 [GPU: attention] → [CPU: MoE 的 up/gate/down] → [GPU: 下一层 attention]。于是第 2、3 层之间会形成两个 GPU split：前一段以第 2 层 attention 链条的第一个节点开头，后一段以 GGML_OP_CPY 开头，那个 CPY 负责把 CPU 上的专家输出搬回 GPU。第 1 层、第 2 层前半、第 2 层后半这三段各有不同的首节点，各自进入一张独立的 ggml_cuda_graph。如果拿 split 序号做键，一旦切分数量变化，序号 1 可能从「第 2 层 attention 段」变成「第 5 层 FFN 段」，图就会被张冠李戴；而按首节点取键，新的计算段自然拿不到旧的捕获对象。

还有一个细节：首节点指针稳定并不等于「同一张 split」——它只说明「这一段计算的入口张量分配在同一个地址」。真正的等价性判定交给后面的 uid/node_props，所以键选得「够稳」就行，不需要精确。

第二道关卡：uid 快捷通道与逐节点比对

1
// 注：为便于阅读，下面把调试日志与空指针判断等语句合并省略了少许行
2
static bool ggml_cuda_graph_update_required(ggml_backend_cuda_context * cuda_ctx,
3
                                            ggml_cgraph * cgraph) {
4
    bool res = false;
5

6
    const void * graph_key = ggml_cuda_graph_get_key(cgraph);
7
    ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key);
8

9
    // ① uid 相同 → 直接返回 false，不做逐节点比对
10
    if (cgraph->uid != 0 && cgraph->uid == graph->uid) {
11
        if (ggml_cuda_graph_debug_enabled()) {
12
            GGML_LOG_DEBUG("CUDA Graph id %zu reused\n", cgraph->uid);
13
        }
14
        GGML_ASSERT((int)graph->node_props.size() == cgraph->n_nodes);
15
        return false;
16
    }
17

18
    graph->uid = cgraph->uid;
19

20
    // ② 节点数变了 → 需要更新
21
    if ((int)graph->node_props.size() != cgraph->n_nodes) {
22
        res = true;
23
        graph->node_props.resize(cgraph->n_nodes);
24
    }
25

26
    // ③ 逐节点比对属性快照
27
    for (int i = 0; i < cgraph->n_nodes; i++) {
28
        ggml_cuda_graph::node_properties prop = {};
29
        memcpy(&prop.node, cgraph->nodes[i], sizeof(ggml_tensor));
30

31
        for (int j = 0; j < GGML_MAX_SRC; ++j) {
32
            if (cgraph->nodes[i]->src[j]) {
33
                prop.node_src_data_ptrs[j] = cgraph->nodes[i]->src[j]->data;
34
                memcpy(prop.node_src_ne[j], cgraph->nodes[i]->src[j]->ne, sizeof(prop.node_src_ne[j]));
35
                memcpy(prop.node_src_nb[j], cgraph->nodes[i]->src[j]->nb, sizeof(prop.node_src_nb[j]));
36
            }
37
        }
38

39
        if (res || memcmp(&graph->node_props[i], &prop, sizeof(prop)) != 0) {
40
            graph->node_props[i] = prop;
41
            res = true;
42
        }
43
    }
44

45
    return res;
46
}

uid 快捷通道从哪来。 sched->splits[i].graph.uid = ggml_graph_next_uid(); 在 ggml_backend_sched_split_graph 里被逐 split 赋值。而 split_graph 只在 sched_alloc_graph/sched_reserve 中被调用。关键问题是：解码时调度器会每步重新 alloc 吗？不会——ggml_backend_sched_graph_compute_async 的判断是：

1
if (!sched->is_reset && !sched->is_alloc) {
2
    ggml_backend_sched_reset(sched);
3
}
4
if (!sched->is_alloc) {
5
    if (!ggml_backend_sched_alloc_graph(sched, graph)) {
6
        return GGML_STATUS_ALLOC_FAILED;
7
    }
8
}
9
return ggml_backend_sched_compute_splits(sched);

只要调度器上一次已经 alloc 过、之后没有被 reset，同一个图结构就在持续复用，split 的 uid 保持不变。于是解码时，绝大多数步骤在 ① 处就短路返回「无需更新」，只做一次 cudaGraphLaunch。

这里的因果链值得写清楚，因为它把「上层图复用」和「下层图回放」连在一起：

• llama_context::process_ubatch 在 can_reuse(gparams) 为真时不重建图，sched 也不会被 reset，下一步 graph_compute_async 直接跳过 alloc_graph，于是 split_graph 不再执行，uid 不变；
• 一旦 can_reuse 为假（序列数、输出数、图参数、记忆上下文任何一项不兼容），res->reset() 与 ggml_backend_sched_reset 都会发生，uid 重新发号，① 不再命中，于是走 ②③ 的逐节点比对；
• 逐节点比对的结果，才决定「是不是真的要重新捕获」的后续判定。

也就是说：上层「图能不能复用」的判定，直接决定了后端「图能不能直接回放」的判定是否要付出比对成本。两者用的是两套不同的比较逻辑（llm_graph_params::allow_reuse 比参数，后端比节点属性），但答案在大多数情况下一致。这也解释了为什么代码里对 can_reuse 有一句谨慎的注释：为了让复用正确，图的完整拓扑必须由这些图参数唯一决定。若这个前提被破坏，后端侧还有一次「逐节点比对」兜底，代价是每步一次全图快照比对。

ggml_backend_sched_synchronize 里还有一段相关注释：图没有处于 allocated 状态时，同步之后总把 next_copy 归零，使得生成阶段每次使用同一个拷贝缓冲，避免图因为在两个拷贝之间来回切换而失效。

顺带一提，ggml_backend_sched_synchronize 里有一段相关注释：图没有处于 allocated 状态时，同步之后总把 next_copy 归零，使得生成阶段每次使用同一个拷贝缓冲，避免图因为在两个拷贝之间来回切换而失效。

漏掉 uid 会让什么出错。 如果某个 split 的 uid 恰好被另一张形状不同的图复用了（比如 pp 图与 tg 图的某个 GPU split 在同一调度器上先后 alloc），① 就会误判。llama.cpp 通过两点规避：一是 split_graph 对每个 split 都重新发号，不同形状的图在 alloc 时得到不同 uid，不会在 ① 处错配；二是即便 ① 通过，② 的节点数断言 GGML_ASSERT((int)graph->node_props.size() == cgraph->n_nodes) 也会在节点数不一致时直接 abort，而不是静默回放一张错的图。

属性快照为何存三样。 三项分别覆盖三类「图里会变的东西」：

这份快照有多大、比对有多贵。 按结构体字段估算：node_properties 里存了一份完整的 ggml_tensor（含 ne/nb、op_params、src[10] 指针等，几百字节），再加上 10 个源的 data 指针（80 字节）、10 组 ne（10×4×8＝320 字节）与 10 组 nb（同样 320 字节），单节点约在 1 KB 量级。一张上千节点的大图，一次全量 memcmp 就是约 1 MB 的内存读。单看它不足以主导一个几十毫秒的解码步，但它与「一次回放省下的提交开销」处在同一量级——若每步都做全量比对，图化的净收益会被抵消掉相当一部分。这就是 uid 快捷通道存在的意义：形态稳定时，比对应当被跳过。

第三项是最容易被忽略的。ggml_tensor 里 ne 表示逻辑形状、nb 表示步长；很多「原地复用的 buffer」在解码时地址不变，但作为某个算子的源时形状含义不同（例如 VIEW 之后 reinterpret 成不同的行数）。如果只比对数据地址，就会漏掉这类差异。

第三道关卡：可捕获性检查

1
static bool ggml_cuda_graph_check_compability(ggml_cgraph * cgraph) {
2
    bool use_cuda_graph = true;
3
    for (int i = 0; i < cgraph->n_nodes; i++) {
4
        ggml_tensor * node = cgraph->nodes[i];
5
        if (ggml_cuda_is_view_or_noop(node)) {
6
            continue;
7
        }
8
        // [TAG_MUL_MAT_ID_CUDA_GRAPHS]
9
        if (node->op == GGML_OP_MUL_MAT_ID) {
10
            const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
11
            if (ggml_cuda_mul_mat_id_needs_sync(node, cc)) {
12
                // the mul_mat_id fallback path synchronizes the stream,
13
                // so we cannot use CUDA graphs
14
                use_cuda_graph = false;
15
                ...
16
            }
17
        }
18
    }
19
    return use_cuda_graph;
20
}

整个检查只排查一种「致命算子」：GGML_OP_MUL_MAT_ID，也就是 MoE 的按专家分组的矩阵乘；同一类算子在别的框架里也有对应实现，比如 vLLM 的 FusedMoE 层与 SGLang 的门控 Router 与按专家对齐。原因是 ggml_cuda_mul_mat_id 有一条 fallback：当类型不是 F32，或本次行数超过 MMVQ/MMQ/MMF 的适用上限时，它会把输入按专家「排序」成一段连续的 get_rows 矩阵乘：ids 的排序索引先在主机上构造出来（std::vector<int32_t> ids_to_sorted_host 那一套），回传 GPU 后还要用 cudaStreamSynchronize 等排序 kernel 完成，再发真正的矩阵乘 kernel。捕获期间不允许这种同步，所以整张 GPU split 直接放弃图化——注意是整张，而不是只有这一个节点退回 eager。

ggml_cuda_mul_mat_id_needs_sync 只做「与内核能力相关的判定」，不做任何形状之外的复杂分析：

1
static bool ggml_cuda_mul_mat_id_needs_sync(const ggml_tensor * dst, const int cc) {
2
    const ggml_tensor * src0 = dst->src[0];
3
    const ggml_tensor * src1 = dst->src[1];
4

5
    if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) return true;
6

7
    if (dst->ne[2] <= MMVQ_MAX_BATCH_SIZE) {
8
        if (ggml_is_quantized(src0->type)) {
9
            if (dst->ne[2] <= get_mmvq_mmid_max_batch(src0->type, cc)) return false;
10
        } else if (GGML_CUDA_CC_IS_AMD(cc)) {
11
            return false;
12
        }
13
    }
14
    if (ggml_cuda_should_use_mmq(src0->type, cc, src1->ne[2], /*n_experts=*/src0->ne[2])) return false;
15
    if (ggml_cuda_should_use_mmf(src0->type, cc, warp_size, src0->ne, src0->nb, src1->ne[2],
16
                                 /*mul_mat_id=*/true)) return false;
17
    return true;
18
}

注意它关注的是 dst->ne[2]，也就是这次 MoE 层要处理的行数（token 数 × 每 token 专家数），而不是「专家数量」。对解码阶段，ne[2] 通常只有几行到几十行，必然能落进 MMVQ 的快速路径，因此不会触发关闭；而预填充大批量时 ne[2] 很大，可能进 MMQ/MMF——只要进了这三条任何一条，就 return false（无需同步）。真正导致关闭的是「三条快速路径都不满足」，常见于非常规 dtype（比如非 F32 的输出）或 AMD 上非量化的特殊组合。

mmq / flash-attn 这些 kernel 的内部细节已在 llama.cpp 源码完全拆解（七）：后端与 Kernel 里拆过，这里只需要记住一个结论：MoE 层的图化与否，取决于 MoE 矩阵乘当时走的是哪条 kernel（MMVQ/MMQ/MMF），而这又与「量化类型 + 硬件 + 本次行数」三者共同决定。对同一份模型、在同一个 GPU 上，解码时基本恒定，所以要么一直能图化，要么一直不能；但每次开关 --n-cpu-moe 或更换量化版本，都可能改变结论。

第四道关卡：热身状态机——为什么必须连续两次稳定

1
static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, ggml_cgraph * cgraph) {
2
    ...
3
    graph_key = ggml_cuda_graph_get_key(cgraph);
4
    ggml_cuda_graph_set_enabled(cuda_ctx, graph_key);
5
    ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key);
6

7
    if (graph->is_enabled()) {
8
        const bool graph_compatible = ggml_cuda_graph_check_compability(cgraph);
9
        if (graph_compatible) {
10
            const bool properties_changed = ggml_cuda_graph_update_required(cuda_ctx, cgraph);
11

12
            if (!graph->warmup_complete) {
13
                // Warmup: need at least 2 calls with no property change on the 2nd call
14
                if (!properties_changed) {
15
                    graph->warmup_complete = true;
16
                    use_cuda_graph = true;
17
                    cuda_graph_update_required = true;
18
                }
19
            } else {
20
                if (properties_changed) {
21
                    graph->warmup_complete = false;   // 属性变了 → 重置热身
22
                } else {
23
                    use_cuda_graph = true;
24
                    cuda_graph_update_required = graph->instances.empty();
25
                }
26
            }
27
        }
28
    }
29
    ...
30
}

这段逻辑可以这样理解。第一次遇到某张图时，update_required 必然返回 true：这个 ggml_cuda_graph 是新对象，uid 为 0、node_props 为空，逐节点比对必然发现「全部需要更新」，于是本次只做 eager 执行、顺便把第一轮属性写进快照。第一次不能直接捕获，因为「这张图第一次被跑」时的状态往往还不是稳态：cuBLAS/cuBLASLt 的句柄与 workspace 可能在这时初始化、部分内核模块可能尚未加载、调度器的 buffer 也可能在这一步才定型。要求「第二次属性仍然一致」才捕获，相当于用一次完整的前向确认「地址与形状已经稳定」，避免把只属于首轮的状态录进图里。

• 第 1 次：update_required=true，不捕获，eager 执行，记下属性；
• 第 2 次：update_required=false → 置 warmup_complete=true，捕获并执行；
• 第 3 次起：update_required=false → cuda_graph_update_required = instances.empty()，此时 instances 非空，直接 cudaGraphLaunch 回放。

如果中途属性变了：warmup_complete 被重置为 false，下一次又回到 eager 执行、重新记快照。这种「变了就重新热身」的代价是：形状频繁变动的场景（比如每步 n_tokens 都不同的 prefill）根本到不了回放——每次都是 eager。这是 llama.cpp 的图化只对解码阶段（batch 稳定、n_tokens 固定）有效的原因。

还要补一个只在上方骨架里省略掉的细节：架构判定写在 ggml_cuda_graph_set_enabled 里，而不是 is_enabled() 中。该函数在 instances 还是空的时候（也就是这张图第一次出现）检查设备计算能力，低于 Volta（compute capability 7.0）就把 disable_due_to_gpu_arch 置真；之后 is_enabled() 会一票否决。把这段逻辑放在「首次出现」时做，是为了避免每步重复查询设备属性，也避免在已有实例之后突然改变策略。

第五道关卡：捕获主循环——融合、并发流、分块

真正的执行与捕获逻辑在 ggml_cuda_graph_evaluate_and_capture。它的外层是这样的：

1
while (!graph_evaluated_or_captured) {
2
    // Only perform the graph execution if CUDA graphs are not enabled, or we are capturing the graph.
3
    if (!use_cuda_graph || cuda_graph_update_required) {
4
        ...
5
        for (int i = 0; i < cgraph->n_nodes; i++) {
6
            ggml_tensor * node = cgraph->nodes[i];
7

8
            // ① 分块边界
9
            if (use_cuda_graph && chunk_size > 0 && i >= chunk_end && !is_concurrent_event_active) {
10
                ggml_cuda_graph_launch_chunk(cuda_ctx, graph, n_chunks++);
11
                CUDA_CHECK(cudaStreamBeginCapture(cuda_ctx->stream(), cudaStreamCaptureModeRelaxed));
12
                chunk_end = i + chunk_size;
13
            }
14

15
            // ② 并发流的 fork / join
16
            if (is_concurrent_event_active) {
17
                if (node == concurrent_event->join_node) {
18
                    ... // 汇合：record join_events + streamWaitEvent
19
                    is_concurrent_event_active = false;
20
                } else {
21
                    cuda_ctx->curr_stream_no = concurrent_event->stream_mapping[node];
22
                }
23
            }
24

25
            prev_i = i;
26

27
            // ③ 视图/no-op 节点直接跳过
28
            if (ggml_cuda_is_view_or_noop(node)) continue;
29

30
            // ④ 非计算标志的节点直接跳过
31
            if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) continue;
32

33
            // ⑤ 融合尝试：返回跳过多少个后续节点
34
            int nodes_to_skip = ggml_cuda_try_fuse(cuda_ctx, cgraph, i);
35
            if (nodes_to_skip != 0) {
36
                i += nodes_to_skip;
37
                continue;
38
            }
39

40
            // ⑥ 普通单节点执行
41
            bool ok = ggml_cuda_compute_forward(*cuda_ctx, node);
42
            GGML_ASSERT(ok);
43
        }
44

45
        if (use_cuda_graph && cuda_graph_update_required) {  // End CUDA graph capture
46
            ggml_cuda_graph_launch_chunk(cuda_ctx, graph, n_chunks++);
47
            // drop the chunks left over from a longer graph
48
            for (size_t i = n_chunks; i < graph->instances.size(); ++i) {
49
                CUDA_CHECK(cudaGraphExecDestroy(graph->instances[i]));
50
            }
51
            graph->instances.resize(n_chunks);
52
            graph_evaluated_or_captured = true;
53
        } else {
54
            graph_evaluated_or_captured = true;
55
        }
56
    }
57

58
    if (use_cuda_graph && !cuda_graph_update_required) {
59
        // the chunks were launched during capture, otherwise launch them here
60
        for (cudaGraphExec_t instance : graph->instances) {
61
            CUDA_CHECK(cudaGraphLaunch(instance, cuda_ctx->stream()));
62
        }
63
    }
64
}

这里面有几个非常值得展开的点。

1
ggml_cuda_concurrent_event concurrent_event(nodes_per_branch.size());
2
concurrent_event.join_node = join_node;
3
for (size_t branch_idx = 0; branch_idx < nodes_per_branch.size(); branch_idx++) {
4
    for (const ggml_tensor * n : nodes_per_branch[branch_idx]) {
5
        concurrent_event.stream_mapping[n] = branch_idx + 1;   // 分支 0 对应流 1
6
    }
7
}
8
...
9
concurrent_events.emplace(root_node, std::move(concurrent_event));

1
static int ggml_cuda_graph_chunk_size() {
2
    static const int chunk_size = [] {
3
        const char * env = getenv("GGML_CUDA_GRAPH_CHUNK_SIZE");
4
        return env != nullptr ? atoi(env) : 512;
5
    }();
6
    return chunk_size;
7
}

1
static void ggml_cuda_graph_launch_chunk(ggml_backend_cuda_context * cuda_ctx,
2
                                         ggml_cuda_graph * graph, size_t idx) {
3
    cudaGraph_t cuda_graph = nullptr;
4
    CUDA_CHECK(cudaStreamEndCapture(cuda_ctx->stream(), &cuda_graph));
5

6
    if (idx < graph->instances.size()) {
7
        ggml_cuda_graph_update_executable(graph->instances[idx], cuda_graph);
8
    } else {
9
        GGML_ASSERT(idx == graph->instances.size());
10
        graph->instances.emplace_back(nullptr);
11
        CUDA_CHECK(cudaGraphInstantiate(&graph->instances[idx], cuda_graph, nullptr, nullptr, 0));
12
    }
13
    CUDA_CHECK(cudaGraphDestroy(cuda_graph));
14

15
    CUDA_CHECK(cudaGraphLaunch(graph->instances[idx], cuda_ctx->stream()));
16
}

1
static void ggml_cuda_graph_update_executable(cudaGraphExec_t & instance, cudaGraph_t cuda_graph) {
2
    cudaGraphExecUpdateResultInfo result_info;
3
    cudaError_t stat = cudaGraphExecUpdate(instance, cuda_graph, &result_info);
4

5
    if (stat == cudaErrorGraphExecUpdateFailure) {
6
        // The pre-existing graph exec cannot be updated due to violated constraints
7
        // so instead clear error and re-instantiate
8
        (void)cudaGetLastError();
9
        CUDA_CHECK(cudaGraphExecDestroy(instance));
10
        instance = nullptr;
11
        CUDA_CHECK(cudaGraphInstantiate(&instance, cuda_graph, nullptr, nullptr, 0));
12
    } else {
13
        GGML_ASSERT(stat == cudaSuccess);
14
    }
15
}

捕获前后与显存池的耦合

ggml_cuda_graph 的生命周期还和显存池绑在一起。CUDA 图里记录的是 buffer 的具体地址，如果这块 GPU 内存被池子回收给别的用途，图再回放就会踩到别人的数据。所以 ggml_cuda_pool 在以下两种情况下都会把所有 cuda_graphs 清空：

1
void invalidate_cuda_graphs() {
2
    int n_captured = 0;
3
    for (const auto & [key, graph] : ctx.cuda_graphs) {
4
        n_captured += graph->instances.size();
5
    }
6
    if (n_captured > 0) {
7
        GGML_LOG_DEBUG(GGML_CUDA_NAME " pool[%d]: memory freed, %d captured CUDA graphs "
8
                      "must be captured again\n", device, n_captured);
9
    }
10
    ctx.cuda_graphs.clear();
11
}

一是 alloc 失败（cudaErrorMemoryAllocation）时先 cudaDeviceSynchronize()、clear_pool()，再 invalidate_cuda_graphs() 后重试分配；二是 free 时池子满了（超过 MAX_BUFFERS）无法缓存、只能 cudaFree。两个场景的共同点是「池子里的地址会发生变化」，因此所有基于 pool 地址录下的图全部作废。注意 clear() 销毁的是 ggml_cuda_graph 对象，其析构函数会把 instances 里每个 cudaGraphExec destroy 掉——地址回收是显存的，但图资源的回收由这套析构负责。

这条规则对推理部署有直接含义：显存紧张时（例如并发请求突增、临时 buffer 扩容），一次池子回收就会让接下来的若干步退回 eager 跑，因为图被全部清空、需要重新经历「热身两次 + 捕获」。这不是 bug，而是「图里地址即身份」这一语义的必然代价。也正因为如此，llama.cpp 在构建上下文时就分别用 pp 与 tg 两张最坏情况图各做一次 reserve，把 buffer 峰值一次性撑到最大，运行期尽量不再 realloc——这段记忆系统与图预留的细节在 llama.cpp 源码完全拆解（四）：llama_context 与记忆系统 里有完整拆解。

图对象按需创建、按时间淘汰

1
ggml_cuda_graph * cuda_graph(const void * first_node_ptr) {
2
    const int64_t time_now = ggml_time_us();
3

4
    // sweep every 5s, evicting cuda graphs unused for >=10s
5
    if (time_now - last_graph_eviction_sweep >= 5'000'000) {
6
        last_graph_eviction_sweep = time_now;
7
        for (auto it = cuda_graphs.begin(); it != cuda_graphs.end(); ) {
8
            if (time_now - it->second->last_used_time >= 10'000'000) {
9
                it = cuda_graphs.erase(it);
10
            } else {
11
                ++it;
12
            }
13
        }
14
    }
15

16
    auto it = cuda_graphs.find(first_node_ptr);
17
    if (it == cuda_graphs.end()) {
18
        it = cuda_graphs.emplace(first_node_ptr, std::make_unique<ggml_cuda_graph>()).first;
19
    }
20
    it->second->last_used_time = time_now;
21
    return it->second.get();
22
}

这是一个很朴素的「LRU 近似」：5 秒扫一次，10 秒没用的图删掉。动机很好理解：llama.cpp 的一个 server 实例同时跑多个 slot、多个序列，不同的 n_tokens 组合可能产生几十种不同的图；这些图里每个 cudaGraphExec 都要占用驱动侧资源（主机内存与内部对象），不做回收会越积越多。10 秒的阈值意味着「短时间内的请求形态波动」不会导致图被销毁，而「长时间换了一种负载」则会被自动清掉，回到冷启动的两步热身。

这里与前面的 last_used_time 字段呼应：捕获、回放、比对，每次都会更新它，从而保证正在用的图不会被回收。

① 断言与 abort 的位置：GGML_ASSERT((int)graph->node_props.size() == cgraph->n_nodes)

在 update_required 的 uid 快捷通道里。如果键命中但节点数不一致，说明「同一个首节点下出现了结构不同的两张图」——在拓扑随配置变化（如设备分配改变、融合开关变化）时理论上可能出现。源码选择 abort 而不是静默处理，因为此时图已被错误地关联到同一个 ggml_cuda_graph，如果继续执行，回放的一定是错的图。这属于「设计上不应该出现，出现就立刻停」的保护。

② cudaGraphExecUpdate 失败后的错误必须清理

(void)cudaGetLastError(); 那句不是多余的。CUDA runtime 返回 cudaErrorGraphExecUpdateFailure 后，错误会留在线程本地的 last error 上；后续任何 CUDA_CHECK 都会读到它并误报。所以重建之前必须清掉。

③ 融合函数返回值的语义不是「融合了多少节点」，而是「主循环应跳过的节点数」

return 5; 表示从当前 i 起共 6 个节点（含 i）由一次融合 kernel 完成，主循环 i += nodes_to_skip 后 ++i，下一个处理的正是 i + 6。写死数字（而不是 sizeof(array) - 1 推导）在这个文件里很常见，因为融合函数的模式是显式列出的；但这使得「模式扩展」时必须同步改返回值。任何一次改错都会导致「少跳一个 → 重复发射」或「多跳一个 → 漏发」，而后者在图里尤其危险：漏发的节点可能在后面仍被当成已产出，结果被后续 kernel 读到未初始化数据。

④ 在捕获模式下，curr_stream_no 与「非 0 流的切换」必须与块边界安全配合

主循环在 is_concurrent_event_active 为真时不允许切块。如果不是这样，块边界落在分支内部就会出现：前一块录入了 fork 与部分分支、后一块直接从「流号已经不在 0」的状态开始录制，导致后一块的录制缺少 fork 事件，回放时它要等到一个从未被 record 的 fork event。

关掉融合来排查，是有效的对照实验

GGML_CUDA_DISABLE_FUSION 常被当作「关融合看看有没有问题」的开关，在排查图化故障时它也确实是一个好用的对照，但要说清楚它能对照出什么、不能对照出什么：

• 能对照的：融合改变了「哪些 ggml 节点会被跳过」，因此改变了实际发射的 kernel 序列，也改变了分块边界落点（切块只在不被融合吸收的节点前触发）。如果「开融合时回放正常、关融合时异常」，优先怀疑块边界与并发事件、以及捕获时 curr_stream_no 的切换位置。
• 不能对照的：执行期融合（try_fuse 里的 MMVQ/MMF 选择、GLU 合并等）不会改变 ggml 图的节点列表，所以它不会改变 node_props 的长度，也不会让属性比对「因为融合不同而判定图变了」。属性比对只看节点、源指针与形状；形状一变（batch 变），它无论如何都会判定「需要更新」。

真正会改变「图拓扑」的，是建图期就能决定的融合（resolve_fused_ops 与模型图里显式生成的融合节点）以及 graph_optimize 给分配器登记的依赖；这部分必须与 batch 无关，否则每次 ubatch 变化都会触发 buffer 重新 reserve，连带把已经捕获的图全部作废。