系列走到第九篇，前两篇把 CUDA 算子优化的基础打法过了一遍——访存模式、bank conflict、指令级并行这些老生常谈的东西，今天换个角度，聊推理场景里真正拉开差距的三个方向：推理 Kernel 的设计思路、CUDA Graphs 的工程落地，以及多 GPU 推理的调度与通信优化。这三个方向有个共同点：它们解决的不是单个算子怎么算得更快，而是整个推理链路怎么不浪费。
如果你正在做模型部署上线，或者想把线上推理服务的吞吐再往上顶一截，这篇文章就是给你准备的。我会把推理 Kernel 和训练 Kernel 的本质区别、融合的层级和取舍、CUDA Graphs 的捕获与重放细节、多卡推理的通信带宽账全部串起来讲，最后附上我在实际项目中踩过的坑和一套可复用的排查清单。内容不绕弯子，全是能直接抄作业的东西。
1. 推理 Kernel 优化：部署场景下的算子设计思路
1.1 训练 Kernel 与推理 Kernel：同源不同命
很多人有一个惯性思维：训练时跑得快的 kernel，推理也一样快。这个想法在单次算子层面没错，但放进整个推理服务里，问题就大了。
训练阶段的数据流是确定性的。batch size 固定，输入 shape 固定，每个 iteration 都是同一套 kernel 序列在跑。你优化的是“单位时间处理多少样本”，也就是吞吐。推理则完全不同，线上请求是动态到达的，batch 可能一会儿是 1，一会儿积压成 64。同一个 kernel，在 batch=1 的场景下，启动开销占比可能高达 30% 以上，这时候你再怎么优化 kernel 内部的访存都没用——瓶颈根本不在计算，而在每次启动的空转。
我踩过最典型的一次坑，是把训练管线里的整套 kernel 直接搬到推理服务里。结果压测时发现 GPU 利用率看着很高，但请求延迟抖动得非常厉害。后来用性能分析工具一看，单个请求的完整链路里有接近 40 次 kernel 启动，加上每次请求之间的设备同步，CPU 侧的开销轻轻松松吃掉几毫秒。这个问题的本质是：训练 Kernel 的设计目标是大 batch 下的高吞吐，而推理 Kernel 需要的是小 batch 下的低开销和可控延迟。
所以做推理优化，第一步不是看 kernel 内部怎么写，而是先把链路里的启动开销、同步点、内存分配全部梳理一遍。这也是为什么后面要讲 CUDA Graphs——它能从机制上消掉大部分启动开销。
1.2 推理场景下 Kernel 融合的三个层级
Kernel 融合不是什么新鲜概念，但推理场景下的融合有自己的层次。我把它们分成三档，按性价比排序。
第一档是算子内部的融合。最典型的就是激活函数与卷积/矩阵乘的融合。比如 Conv+Bias+ReLU 这种结构，训练时你可以分开跑，但推理时完全没必要让中间结果写回显存再读出来。业界最常见的做法是直接用 cudnn 提供的 fused 接口，或者在自定义 kernel 里把 bias 加法和激活函数直接算在输出 stage 中。这一档融合能省掉的是一次全量显存读写——对于大特征图来说，这个量非常可观。
第二档是相邻算子的启动合并。这个不需要改算子逻辑，只需要把多个 kernel 的 launch 放到同一个 CUDA stream 里，让它们按序执行。GPU 本身可以异步执行，只要 CPU 侧持续发射 kernel，就能掩盖一部分启动开销。但这里有个前提：kernel 之间不能有隐式的同步点。很多人写着写着就调用了 cudaDeviceSynchronize，这就把整个流水线打成了串行，再多的 launch 合并都是白费。
第三档是跨层融合，比如把 LayerNorm 和前面的矩阵乘融合，或者把 Attention 中的 QKV 投影合并成一次大 GEMM。这一档收益最大，但实现难度也最高。因为它往往需要同时改动模型结构和 kernel 实现，而且在动态 shape 场景下，融合 kernel 的边界条件处理非常繁琐。我的建议是，先做第一档和第二档，把最容易拿到的收益拿到手，再根据 profile 结果决定要不要上第三档。
1.3 内存分配策略：推理服务里看不见的杀手
推理服务里最常见的性能杀手不是算子本身，而是内存分配。cudaMalloc 是一个同步操作，它的开销远高于一般的 CPU 内存分配，而且高频调用的时候还会引起显存碎片化。训练场景无所谓，反正显存是预先规划好的；推理场景不一样，请求动态到达，如果你在每次请求的路径里都调用 cudaMalloc，那整个服务的延迟会非常难看。
正确的做法是显存池化。在服务启动阶段就把推理需要的内存一次性分配好，后续请求从池里复用。主流推理框架内部都有缓存分配器，原理都一样：先按模型结构预分配一大块显存，然后根据 tensor shape 做 sub-allocation，用完还回池里。这样做的好处不仅是省掉了 malloc 的开销，还能让显存使用保持稳定，不会出现跑了一段时间之后显存碎片化导致分配失败的问题。
另外还有一个容易被忽略的点：锁页内存。Host 侧和张量传输的数据如果放在可换页内存里，cudaMemcpy 会先做一次隐式的页面锁定，这个操作耗时且不可控。推理服务里如果频繁做 host 到 device 的数据传输，一定要用 cudaMallocHost 分配锁页内存。我们实际测下来，在小 batch 场景下这个优化能把单次拷贝的耗时降低 50% 以上，因为省掉了页锁定的开销。
2. CUDA Graphs：把 CPU 从启动马拉松里解放出来
2.1 CUDA Graphs 的工作原理与捕获模式选择
CUDA Graphs 的出现，本质上是为了解决 kernel 启动开销的问题。常规情况下，CPU 提交一个 kernel 到 GPU 执行，中间涉及 launch 命令的解析、参数校验、依赖检查、硬件队列提交等一系列操作，单次开销在 3 到 10 微秒。一个 Transformer 模型前向推理链路上有几十个 kernel，累计起来就是几百微秒的纯 CPU 开销。注意，这个开销不是 GPU 在算，而是 GPU 在等 CPU 下命令，完全是空转。
CUDA Graphs 的思路很直接：把整条 inference 的 kernel 依赖图捕获下来，实例化成一份可重用的图对象，之后每次推理只需要一次 cudaGraphLaunch 就能把整条流水线提交给 GPU。启动开销从“kernel 数量 × 单次启动开销”变成了一次固定的图提交开销。
捕获有三种方式。第一种是流捕获（Stream Capture），这是最常用的方式。你只需要在非默认流上调用 cudaStreamBeginCapture，然后像平常一样 launch 所有 kernel，最后调用 cudaStreamEndCapture，就能拿到一张图。第二种是用 API 手动构建图节点，灵活度最高，但代码量很大。第三种是基于现有图做修改和更新，适用于动态 shape 场景。
流捕获的代码长这样：
cudaStream_t stream; cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal); // 这一整段会被捕获进图里 kernel_attention<<<grid, block, smem, stream>>>(...); kernel_mlp<<<grid, block, smem, stream>>>(...); kernel_norm<<<grid, block, smem, stream>>>(...); cudaStreamEndCapture(stream, &graph); cudaGraphExec_t instance; cudaGraphInstantiate(&instance, graph, NULL, NULL, 0); // 之后每次推理只需要一行 cudaGraphLaunch(instance, stream);这里有一个非常重要的细节：捕获期间 kernel launch 使用哪个 stream，后续的 graph 执行也用哪个 stream。如果你在捕获过程中调用了 cudaDeviceSynchronize 或者让默认流参与了执行，捕获往往会失败或者产生未定义行为。原因在于默认流和其他流之间有隐式的同步关系，这会破坏图捕获的依赖分析。
2.2 动态 Shape 场景下的图管理与多图策略
CUDA Graphs 最大的短板在于它捕获的是静态依赖关系和静态参数。如果你的模型输入 shape 会动态变化，比如 batch size 在 1 到 32 之间变化，那同一张图就不能复用到所有输入上。
工程上有两种解法。第一种是固定 batch size，把请求排队成固定的 batch。服务端配置几个固定的 batch 档位，比如 1、4、8、16、32，每个档位捕获一张图。请求进来之后根据当前排队数量选择最合适的档位执行。这种做法最简单，代价是 batch 大小不匹配时会有一定的吞吐损失。
第二种是使用 cudaGraphExecKernelNodeSetParams 更新图的节点参数。这种方式允许你保留图的结构，只更新每个节点里的 kernel 参数，比如 grid 尺寸或者指针地址。它的开销比重新捕获一张图小得多，但要求模型的 layer 结构完全不变，只有 shape 相关参数变化。实际操作中，这个更新过程仍然有 CPU 开销，所以适合 shape 变化不频繁的场景。
我在实践中更推荐第一种多图策略。原因很简单：推理服务的请求分布通常有一定的规律，固定档位可以做到接近最优的利用率，而且多图管理在代码层面更直观，出问题也好排查。图更新机制适合比较进阶的团队去玩，收益有限，复杂度却不小。
2.3 实测收益与适用边界
在我们的实测中，对一个中小规模的 Transformer 模型做单卡推理，开启 CUDA Graphs 之后，端到端延迟降低了大约 15% 到 30%。这个收益主要来自两方面：一是启动开销的消除，二是整个图执行过程中 CPU 和 GPU 的重叠做得更好，GPU 几乎不需要等待 CPU。
不过 CUDA Graphs 不是万灵药。它有一个比较麻烦的工程问题：捕获期间分配给 kernel 的临时显存地址会被固化进图里，如果这些内存在后续过程中被其他操作复用或释放，图重放时就会出现非法地址访问。因此，如果你用的是自定义内存池，必须确保图捕获期间从池里拿到的内存在图执行期间不会被偷走。
此外，如果模型本身是动态控制流的，也就是 kernel 执行路径取决于 tensor 数据的内容，那 CUDA Graphs 就无法直接捕获，需要预先展开所有分支或者在分支处拆多张图。这个限制在做推理时很致命，因为部署方往往需要保留模型的动态逻辑。我的建议是，在模型导出阶段就把动态控制流尽量转成静态的 shape 计算，这样后期的图上手难度会小很多。
3. 多 GPU 推理：从单卡到多卡的性能工程
3.1 数据并行、张量并行、流水线并行的选择逻辑
多 GPU 推理最常见的动机是：单卡放不下模型，或者单卡吞吐满足不了服务需求。这两个动机会导向完全不同的架构选择。
如果模型能放进单卡，只是吞吐不够，那数据并行是最直接的选择。把多张卡各自部署一份完整的模型，请求进来后按负载均衡策略分发到不同的卡上。这个方案不需要考虑跨卡通信，扩展性最好，是大多数推理服务的第一选择。
如果模型单卡放不下，就必须做模型并行。这里又分两种思路。张量并行是把单个层的矩阵切成多份，分布在多张卡上协同计算，每一层都需要一次跨卡通信。流水线并行则是把模型的层切成多个阶段，每个阶段放在一张卡上，数据按顺序流过各阶段。张量并行延迟低，但通信开销大；流水线并行通信少，但会有流水线气泡，也就是阶段之间的空等。
选择时的判断依据其实就两条。第一，模型参数量除以单卡显存，如果超出不多，优先考虑流水线并行，把通信降到最低。第二，如果模型单层特别大，比如超大维度的 embedding 或者超大 hidden size，那流水线并行也放不下，只能上张量并行。
推理场景还有一个特有因素需要考虑：batch size。张量并行下，每个 batch 都要做跨卡通信，通信延迟和 batch 大小成正比。数据并行则完全不需要通信，但每张卡各自独立处理请求，无法享受大 batch 带来的计算效率提升。所以有的团队会做“混合并行”：数据并行扩容，每路内部再用张量并行或流水线并行处理超大的模型。这种组合拳在超大模型推理服务里几乎是标配。
3.2 通信拓扑与带宽账：多卡推理的关键计算
多 GPU 推理里，通信开销是躲不开的。你首先要知道你的多卡拓扑是 NVLink 还是 PCIe，这决定了通信带宽的上限。NVLink 的带宽通常远高于 PCIe，而且 NVLink 支持 GPU 之间的直接访问，不需要经过 CPU 内存中转。如果条件允许，多卡推理场景优先选支持 NVLink 的硬件平台。
然后我们来算一笔账。以张量并行中的一个 Transformer 层为例，前向计算中每一层 Transformer 都需要做至少两次 AllReduce 通信，一次在 Attention 之后，一次在 MLP 之后。假设模型的 hidden size 是 H，那么每一次 AllReduce 传输的数据量大约为 2×H×batch_size×2 字节（fp16 下）。如果 H=8192，batch=8，一次 AllReduce 的数据量就是 256KB 左右。在 NVLink 上这不算什么，但在 PCIe 上就能明显感知到延迟。
通信时间对整个推理延迟的占比，可以直接用通信数据量除以有效带宽来估算。比如 PCIe Gen4 的单向带宽大约 32GB/s，256KB 的传输也就是几微秒，但你要是加上同步开销和通信库内部的对齐损失，实际耗时可能到 20 到 30 微秒。如果模型有几十层，每一层叠加下来，通信时间就非常可观了。
所以多卡推理的优化核心是让通信和计算尽可能重叠。常规做法是：在 kernel 计算当前层的同时，让通信引擎去传输下一层需要的数据。这需要你手动切分 stream，把 AllReduce 的通信调用放到单独的通信 stream 上，并且用 CUDA Event 来管理依赖关系。没有做重叠的实现，多卡推理的延迟会比单卡慢一个数量级，这不是玩笑。
3.3 多 GPU 推理落地的工程细节
多 GPU 推理的工程实现里，有几个细节特别容易出问题。
第一个是通信库的初始化。使用主流的集合通信库时，每个进程需要初始化独立的通信域，并且进程之间的 rank 映射必须和 GPU 的实际物理拓扑匹配。如果 rank 映射和 GPU 绑定错了，通信会走最慢的路径，性能一落千丈。我们之前踩过这个坑：四卡机器上，把 rank 0 和 rank 1 分配到了同一块 GPU 上，结果通信直接退化成走 CPU 内存的链路，吞吐掉了 70%。
第二个是负载均衡。数据并行部署时，如果请求分发策略做得不好，很容易出现某张卡已经排队排队，另一张卡还在空等。解决方式是使用全局的动态批处理队列，而不是每张卡独立接请求。请求先进入一个全局窗口，窗口满了以后统一按策略分配到各张卡。这样既保证了 batch 的充分性，又避免了单卡过载。
第三个是通信与计算资源的隔离。GPU 上的通信操作和计算操作如果混在同一个 stream 里，可能会导致通信阻塞计算。稳妥的做法是给通信单独分配 stream，并通过事件同步保证计算和通信之间的依赖关系正确。通信 stream 和计算 stream 可以并行运行，但要在需要数据的位置做精确的等待。
多卡推理还有一个必须重视的问题是容错。单卡坏了，整个服务能降级；多卡场景下，某张卡的通信超时或者显存泄漏，会直接影响整个通信域。所以即使只是做推理部署，也建议在启动时做一次全量健康检查，包括通信带宽测试和显存读写测试，避免服务上线后才发现某张卡是亚健康状态。
4. 常见问题与排查技巧：实测里那些让人抓狂的坑
4.1 性能数据不会说谎，但基准测试会
做推理优化的人都经历过一种困惑：明明报告上 GPU 利用率 90% 多，但线上延迟就是下不来。这种情况十有八九是基准测试的方法出了问题。
最常见的误区是直接用训练时的测速脚本去测推理性能。训练脚本通常是连续喂数据，没有真正的请求间隔，测出来的是纯计算吞吐。推理服务不一样，请求是稀疏到达的，中间有空窗期，GPU 在空窗期里就是空转。你测出来的平均利用率高，可能是因为你压测时的请求频率已经超过了服务能处理的上限，请求在排队，GPU 一直在忙，但单请求延迟已经恶化。
正确的方法是分场景测。吞吐场景用持续压测，看稳态吞吐和延迟分布；延迟场景用低并发压测，看 P50、P95、P99 延迟。两个指标的优化方向完全不同，前者关注 GPU 利用率、batch 大小、kernel 效率；后者关注启动开销、同步点、内存分配、通信延迟。混在一起测，你只会得到一团糊涂的数据。
另外一个常见的坑是不做 warm-up。CUDA 在第一次运行 kernel 时会有 JIT 编译和显存初始化的开销，如果不先跑几轮 warm-up，直接测延迟，得到的数据虚高很多。我之前测 CUDA Graphs 收益时，就是因为没做 warm-up，把一部分 JIT 编译开销算到了 baseline 的头上，导致图加速效果被严重高估。
4.2 多卡场景下的调试方法与工具
多卡推理出问题时，最怕的就是“看起来每张卡都正常，但整体性能稀烂”。这种情况需要系统性的排查。
第一步先确认通信链路。用通信库自带的 bandwidthTest 工具测试两两 GPU 之间的实际带宽，和理论值对比。如果明显偏低，检查一下是否是 PCIe 链路降速，或者通信库里的 rank 映射出了问题。这一步能排除绝大多数通信环境的问题。
第二步做分段 profile。不要直接对整个推理链路做 profile，那样数据太多，看不清楚。把链路切成几段：输入预处理、图执行、后处理、通信。分段 profile 之后，你就能定位性能瓶颈到底在哪个阶段。我习惯用性能分析工具抓完整的 kernel 时间线，然后对比不同阶段的时间占比，这比单纯看 GPU 利用率靠谱得多。
第三步是检查同步行为。多卡推理里最隐蔽的性能问题是隐式同步。比如你在两张卡上分别执行了不同的任务，但没有正确地用 event 建立依赖，后续某个操作会把整个设备强制同步一次，导致之前积累的异步执行全部被打断。性能分析工具的时间线里，如果出现大段的空闲间隙，多半就是隐式同步在作祟。
4.3 一套可复用的推理链路排查清单
最后，分享一份我在项目里沉淀下来的排查清单。遇到推理性能问题时，按顺序做一遍，能解决八成问题：
- 先确认瓶颈在 CPU 还是 GPU。用性能分析工具看 GPU 空闲占比，如果 GPU 空闲时间占比高，说明 CPU 发射速度跟不上，优先优化启动开销或合并 kernel。
- 检查是否有隐式同步。搜索代码里的 Synchronize 调用，包括设备级同步和流级同步，逐一确认是否必要。
- 检查内存分配频次。统计运行路径里 cudaMalloc 和 cudaFree 的调用次数，超过两位数就得考虑内存池。
- 检查数据传输路径。确认 host 到 device 的拷贝是否用了锁页内存，是否采用了异步拷贝。
- 多卡场景下，先测单卡性能，再叠加多卡。单卡性能没过关之前，别急着调多卡通信。
- 每次只改一个变量，改完重新压测。推理优化最怕同时改多个配置，最后出了问题都不知道是哪一步引入的。
这份清单不是什么高深理论，但它能帮你快速缩小问题范围，避免在无关的方向上浪费大量时间。
这次先写到这里。CUDA 算子优化这个系列走到今天是第三篇，后续如果有机会，我想专门展开聊聊推理服务里的动态批处理与调度策略——那一块和今天讲的 Kernel 优化、CUDA Graphs 结合起来，才是线上推理性能优化的完整拼图。如果你在实际部署中也碰到过类似的坑，欢迎按这份清单试一试，排掉一两个隐性问题之后再看性能数据，你会有种豁然开朗的感觉。
