企业官网建设流程全解析

做深度学习Runtime的同学，大概都遇到过这样一个尴尬局面：模型移植到自己写的调度器上，端到端跑一遍，GPU利用率低得吓人；打开时间线一看，拷贝引擎干活的时候计算单元在睡大觉，计算单元忙起来的时候DMA又闲着。这个自研项目的主题，就是在Runtime内部把硬件调度和执行流驱动机制真正打通，把所有能并发的任务，用执行流（stream）这种手段铺到GPU的各个处理单元上。简单说，它解决的不是“某个kernel怎么写得快”，而是“几百个kernel在一张卡上怎么排队、怎么穿插、怎么把带宽和算力同时喂饱”。

这个内容适合三种人看：一是做推理引擎开发的，二是做训练框架底层支持的，三是对GPU编程有一定基础、想真正搞懂多流并发原理的人。我尽量把调度器的设计思路、执行流的用法、以及踩过的坑都摊开写，读完你至少能知道，怎么让自己的Runtime在同样的硬件上多跑出一截性能。

1. 从问题出发：为什么Runtime需要一张独立的“调度大脑”

1.1 GPU利用率低，锅往往不在算子而在调度

我们组一开始的做法跟大多数团队差不多：拿到一个模型，先把计算图优化、算子融合做一遍，然后把每个kernel都优化到单测性能看起来很漂亮。结果整张图端到端跑起来，GPU利用率只有百分之四十几。单独测每个算子，个个都快；合到一起就慢，问题出在衔接上——算子在stream里的排放顺序不合理、等待条件太多、host端反复同步。这类调度问题在单算子评测里完全看不见，只有端到端时间线会暴露。

可以简单算一笔账：假设一个模型有60个算子，每个kernel在GPU上真正执行50微秒，理想情况总耗时也就3毫秒。但如果算子之间因为同步和launch排队出现30微秒的空当，60个算子就会多出1.8毫秒的等待，总耗时变成4.8毫秒，利用率直接掉到六成以下。这种损失不是靠优化某个kernel能救回来的，因为你算的每一笔账都会被调度层的漏洞漏掉。

所以这个项目从一开始就把“硬件调度”当成一个独立模块来设计，而不是顺带处理的事。它的职责很简单：决定每个任务什么时候提交、提交到哪条流、要等哪些事件、用什么优先级。把这些定清楚，才能真正压出硬件的并发能力。

1.2 这个项目要解决的三个层面的并发问题

GPU的并发不是单靠开一堆线程就能得到的，它至少分成三个层面：

第一，计算与拷贝的重叠。DMA搬运数据的时候，SM应该在算别的任务，而不是空转。这是最容易见效的一层，也是很多Runtime忽视的一层。

第二，多个kernel在GPU上的并发驻留。一个kernel如果只占用了少量SM，另一个流上的kernel可以同时上卡，充分利用空闲的SM。很多人的多流并发没效果，就是因为流里的kernel太大，一个就霸占了整张卡，其他流只能排队。

第三，多个请求或流水线之间的交错。在线推理场景里，请求是源源不断进来的，不同的请求之间没有强的依赖关系，理论上可以完全并行地推进。但大多数Runtime实现是串行处理完一个请求再处理下一个，白白浪费了并发机会。

项目目标就是把这三层并发统一在一个调度框架里：上层图IR负责静态分析，调度内核负责运行时决策，执行后端把任务真正铺到多流上。这三个层面缺一个，并发性能都上不去。

2. 硬件执行模型：先看懂GPU上的“并发资源”再谈调度

2.1 GPU上真正能并行的是哪些部件

要理解调度，先得知道GPU上到底有哪些可以并行的硬件资源。粗略分四类：

• 计算引擎：由一堆SM组成，写kernel最终都是算在它头上。每个SM能驻留的线程块数量有限，寄存器、共享内存也有限，所以一个kernel不可能无限占满所有SM；同理，多个小kernel可以同时挤在不同SM上。
• 拷贝引擎：负责主机与设备之间、以及设备内存内部的数据搬运。它是独立硬件，可以和计算引擎同时干活，前提是你给的搬运指令本身是异步的。
• host端launch通道：CPU向GPU提交kernel和拷贝指令也需要时间，而且不是零成本。一次launch几微秒到十几微秒，对小kernel来说这个开销可能比内核执行本身还大。
• 显存带宽和互联带宽：所有计算、搬运共享的受限资源，很多算子不是被算力卡住，而是被带宽卡住。

用生活类比来说，GPU就像一个车间：SM是工位上的加工工人，DMA是运料叉车，PCIe/显存带宽是车间通道，host端launch是车间主任往工位派单。如果只有一条流水线，工人加工完所有零件，叉车才去运下一批，那么工人和叉车永远有一个在等。聪明的做法是把料分成几批，叉车在送第i+1批的时候，工人正好在加工第i批。

2.2 执行流（Stream）到底是一种什么语义

一个stream就是一条提交到GPU的任务队列。同一条队列里的任务严格按照提交顺序执行，前一个没干完，后一个不会开始。不同队列之间没有顺序保证，谁先执行完都可能，甚至两个流上的kernel同时上卡执行也完全正常。

默认stream是个大坑。如果不指定stream，所有launch都进默认队列。旧式默认队列有个“独特”行为：它既会等之前非默认流上的任务，又会让后续非默认流等它，相当于一个隐式的全同步栅栏。很多人开了多流并发却发现kernel全部串行，第一嫌疑就是默认流的隐式同步在捣乱。解决办法是创建非阻塞流，或者干脆用per-thread默认流，避免参与隐式同步。

流的并发语义本身很简单，复杂的是让不同流之间按依赖顺序执行。任务B依赖任务A的结果，但A和B在不同流上，这时候必须用事件：在A所在流上放一个“完成标记”，让B所在流等待这个标记。事件就是GPU侧用来表达流间依赖的唯一桥梁。

2.3 并发上限：为什么“Stream越多越好”是最大的错觉

有人一听说多流并发，立刻把每个算子都放到单独一条新流里，结果性能不仅没提升，反而更差了。原因有两个：

第一，GPU上能同时驻留的kernel数量有限。如果第一个kernel就试图占满所有SM，第二个kernel在硬件层面根本插不进来，哪怕在不同的流上也白搭。开了16条流，最后还是一条条的排队。

第二，每条流本身有CPU端维护成本。小kernel只有10微秒的执行时间，但往流里提交一次可能就要花5微秒甚至更多。提交100个小kernel到100条流上，CPU光launch就能把GPU喂吐。正确做法是让kernel粒度适中，流数量匹配任务形态。我后来做推理引擎时，计算流常年维持在4到8条，拷贝流单独1到2条，效果远好于无脑开几十条。

3. 执行流驱动机制：把任务编排成硬件爱吃的形状

3.1 流水线模型：用双缓冲把搬运和计算重叠起来

让计算和拷贝真正重叠，核心思路是“分块流水线”。一次推理的数据搬运不要一次性搬完，而是分成若干组，让拷贝流和计算流交替干活。下面是一个非常典型的双缓冲流水线结构：

// 输入拷贝流、计算流、输出拷贝流各司其职 cudaStream_t stream_copy_in, stream_compute, stream_copy_out; cudaEvent_t ev_in_done[MAX_CHUNK], ev_compute_done[MAX_CHUNK]; for (int i = 0; i < num_chunks; i++) { // 第1步：异步拷贝第i块输入到显存 cudaMemcpyAsync(buf_in[i], host_in + i * chunk_bytes, chunk_bytes, cudaMemcpyHostToDevice, stream_copy_in); cudaEventRecord(ev_in_done[i], stream_copy_in); // 第2步：计算流等待第i块拷贝完成，然后开始算 cudaStreamWaitEvent(stream_compute, ev_in_done[i]); preprocess_kernel<<<grid, block, 0, stream_compute>>>(buf_in[i]); main_kernel<<<grid, block, 0, stream_compute>>>(buf_in[i]); cudaEventRecord(ev_compute_done[i], stream_compute); // 第3步：输出拷贝流等待计算完成，再回传结果 cudaStreamWaitEvent(stream_copy_out, ev_compute_done[i]); cudaMemcpyAsync(host_out + i * chunk_bytes, buf_out[i], chunk_bytes, cudaMemcpyDeviceToHost, stream_copy_out); }

这段代码的关键不是API本身，而是它形成的时序：计算流在处理第i块的时候，输入拷贝流已经在搬运第i+1块；输出拷贝流则在慢慢回收第i-1块的结果。三个引擎各忙各的，谁也不等谁。

这里最容易被忽视的前提是pinned memory。cudaMemcpyAsync要真正异步执行，源和目的地址必须是锁页内存。如果用的是普通可分页内存，API内部大概率会退化成同步拷贝，流水线形同虚设。我在项目里吃过这个亏，后面凡是host端参与拷贝的缓冲全部用锁页内存分配，性能立刻不一样。

3.2 事件依赖的正确使用方法

cudaEventRecord不是让CPU等，它只是把一个里程碑插入到流的执行序列里。真正产生等待效果的是cudaStreamWaitEvent，而且这个调用本身也是异步的——它只是往另一条流里塞了一个“等到某事件后的再继续”的标签，CPU调用完立即返回，GPU侧才排队等待。

常见错误是kernel launch完马上cudaDeviceSynchronize，把整条流水线拦腰截断。正确做法是：如果结果要被另一条流消费，就用事件传递；如果结果最终要被CPU拿回，尽量也用异步拷贝加事件，让最后一个块算完时才真正让CPU等一次。

事件对象的创建和销毁也不便宜，生产环境里一定要做事件池。预分配32个或64个事件，按需取用，等它被消费完再回收。复用事件有一个原则要守住：一个事件刚被record，还没被其他流的wait消费掉，就不要立刻二次record。否则之前设置的依赖关系会错乱，排查起来极其痛苦。

如果一条流要等多个上游事件，就多次调用cudaStreamWaitEvent，GPU侧会等所有里程碑都过了再放行。反过来，一个事件被多条下游流同时等待，也完全合法，这是事件模型的高明之处。

3.3 流优先级和资源分配策略

流优先级是硬件调度里容易被忽略的调节阀。GPU驱动会为高优先级流里的任务分配更多执行机会，适合放延迟敏感的算子；低优先级流放后台数据预取、日志、非关键路径上的杂活。

需要注意两点。一是优先级的方向别记反：先用接口查一次设备支持的优先级范围，数值越小优先级越高，高优先级通常是负数。二是不建议把优先级差拉得太大，否则低优先级流可能被饿死。我见过一个案例，高优先级流里永远有任务在跑，低优先级里的显存拷贝一直排不进去，最终表现为偶发的超时。

流的使用也要池化。每个进程不要频繁创建销毁流，统一维护一个流池，按请求分配、用完归还。尤其注意：不同请求尽量不要塞进同一条流，因为同一条流内是严格串行的，两个并发请求挤一条流会在GPU侧互相排队，等于并发失效。

4. Runtime调度器实现：从任务图到执行流

4.1 总体分层架构

调度器不能跟硬件API揉成一团，否则单元测试很难写。我们最终定下来的分层是这样的：

• 图IR层：承接上层传入的计算图，做算子融合、内存规划、静态依赖分析。
• 调度内核：做拓扑排序，维护就绪队列，为每个任务选择stream、绑定事件依赖。
• 执行后端：负责持有stream池、事件池、显存池，并调用真实的launch接口。
• 分析模块：记录每个任务的提交时间和完成时间，导出时间线，方便和剖析工具对照。

设计意图很简单：调度内核不直接碰硬件API，所有硬件调用都通过执行后端。这样做的好处是，调度内核可以用一个“假启动”的执行后端在CPU上先验证依赖关系，确保没有死锁、没有顺序错乱，再真正上GPU跑。

4.2 依赖编排与就绪队列：一个可持续运行的调度循环

调度器核心是一个事件驱动的状态机。每个任务携带三样东西：它要等哪些事件、自己被分配到哪条流、自己完成后要record哪个事件。调度循环不断从就绪队列里取任务，做完launch，再释放后继任务。

struct Task { int id; int stream_id; vector<cudaEvent_t> wait_events; cudaEvent_t done_event; function<void()> launch_body; // 封装kernel launch }; void ScheduleLoop(vector<Task>& dag, vector<int> dependents[]) { queue<Task*> ready; int remain[MAX_TASKS] = {0}; // 剩余前驱数量 for (int i = 0; i < dag.size(); i++) { if (remain[i] == 0) ready.push(&dag[i]); } while (!ready.empty()) { Task* t = ready.front(); ready.pop(); for (cudaEvent_t ev : t->wait_events) { cudaStreamWaitEvent(streams[t->stream_id], ev); } t->launch_body(); cudaEventRecord(t->done_event, streams[t->stream_id]); for (int nxt : dependents[t->id]) { if (--remain[nxt] == 0) ready.push(&dag[nxt]); } } }

这里最关键的一点是，表达依赖用的是事件而不是“等前驱算子彻底结束”。事件的粒度比kernel粒度细，可以精准表达“哪一步算完了、后续可以继续”的语义，而且不会阻塞host线程。

调度线程本身不能因为GPU执行慢而卡住。实际操作中，调度线程负责依赖逻辑，真正往GPU提交任务可以交给一个launch线程池，让CPU侧launch开销和GPU执行重叠。但有一个大坑：同一流内的任务顺序必须保持提交顺序，多个launch线程同时往同一条流里塞任务可能会乱序。解决办法是在调度阶段就固定每个task的stream归属，并保证所有依赖该流的前置事件都绑定完成后再进入launch队列。

4.3 从Launch到Graph：启动开销的极致压缩

一个推理模型可能有几百个算子，每个算子launch的CPU开销几微秒到十几微秒，加起来就是毫秒级的纯启动损耗。这正是CUDA Graph的用武之地：把固定的一块执行序列捕获成图，提交一次，GPU侧按图内依赖自动调度，CPU侧不再逐个launch。

使用CUDA Graph有前提：模型的shape、分支行为、算子顺序必须是固定的。捕获过程中也不能有阻塞等待和依赖GPU结果的host判断。沾上任何一个，捕获就会失败。动态shape场景怎么办？两种常用手段：一是按shape分组建多个graph，运行时选一个用；二是用节点参数更新接口替换kernel参数，不用重建整张图。

多流并发照样可以捕获进graph。在捕获流上正常使用多个stream配合事件同步，生成的graph会保留流之间的并发关系。效果比单条流捕获出来的图更好，因为graph内部本身就带并行度。这也是我们在优化后期把多流流水线和CUDA Graph结合使用的原因，性能曲线又向上跳了一截。

5. 实操调优：一次多Stream并发优化的完整记录

5.1 从时间线找瓶颈：三步定位法

拿到一个性能问题，我一般先开带时间线的GPU剖析工具，按这三步走：

第一步，把时间线按引擎着色，看拷贝引擎和计算引擎是不是交替空闲。如果看到拷贝引擎干活时计算引擎是空的，说明没有搭起流水线。

第二步，找kernel之间的空白段。空白集中出现在某条流上，说明同步太频繁；空白分散且短，CPU launch排队的原因更大。

第三步，切到host线程视图。如果GPU没活干而CPU忙得像陀螺，基本可以断定launch和同步是瓶颈，而不是算子本身。

这三步做完，问题在调度层还是在计算层就很清楚了。大部分多流并发调优都卡在前两步，真正需要到kernel级剖析的反而少。

5.2 常见问题速查表

表格里每一行都是真实踩过的。特别是默认stream这条，新手最容易中招，而且现象非常隐蔽：代码看起来是多流，时间线里却一条条串着跑，不盯时间线根本意识不到问题。

5.3 实测成绩与参数经验

初期版本是单流、整批拷贝完再计算、再整批拷回，端到端96毫秒，GPU利用率43%。第一次改造用pinned memory加双缓冲、把搬运和计算分成两条流，端到端降到72毫秒，利用率63%。第二次改造把请求按4加4拆成两个计算流，大算子拆块，同时引入事件池和流池，端到端57毫秒，利用率74%。第三次改造把重复执行的子图捕获成CUDA Graph，端到端稳定在48毫秒左右，利用率81%。

这些数字是当时项目里的实测，换张卡换个模型不一定相同，但趋势是一致的：调度层的收益往往比单个kernel优化还大。经验参数上，计算流4到8条、拷贝流1到2条、流水线段数不要少于三段（输入拷贝、计算、输出拷贝）、事件池预分配64个左右，是比较稳妥的起点。具体值要根据任务形态调，但方向不会变。

最后聊点实际体会

项目做完回头看，我最深的体会是：调度层是GPU性能的放大器，但也最容易给人虚假的安全感。很多坑不在代码逻辑上，而在于“看起来异步、其实同步”的API行为。比如我一度以为用了异步拷贝就万事大吉，忘了检查缓冲区是不是锁页内存，结果流水线一直没建立，时间线里看得一清二楚。后来把所有同步点全部过了一遍，性能立刻上了一个台阶。

最后分享一个小技巧：改完调度逻辑先别急着看端到端延迟，先看时间线。如果时间线上有一段GPU空窗而host忙，多半是同步点太多；如果GPU全程在跑但显存带宽接近饱和，那才是算子层面的问题。分清这两个方向，能省掉至少一半的调优时间。这个内容后续还可以往多设备、多卡的方向扩展，不过那就是另一个维度的故事了。