CUDA流与异步操作
说到多GPU负载均衡，有个东西绕不开——CUDA流。我刚开始做 GPU 编程时，总觉得流是个概念货，直到在真实集群里跑 8 卡模型，发现显存复制完全串行，GPU 利用率死活上不去…… 嗯，那之后我才意识到，流是让 GPU 真正“忙起来”的关键。
核心观点：CUDA流本质上是一个异步操作队列。你把一批任务（内核、拷贝、事件等）提交到流里，它们在 GPU 上依次执行，但主机端可以不等它们完成就继续干别的事。你完全可以把它理解为 GPU 版的“任务管道”。
流的本质：队列 + 异步
每个流内部是串行的。操作 A 没做完，操作 B 不会开始。但不同流之间的操作可以并发执行——只要硬件资源允许。
举个例子：你有两张卡，每张卡都要拷贝一批数据，然后跑两个内核。如果全部用默认流，你只能等第一张卡全部搞完，再搞第二张。但如果你为每张卡各创建两个流，数据拷贝和内核计算就能重叠起来。我个人习惯在需要高吞吐的推理服务里，至少创建 4~6 条流。
默认流 vs 非阻塞流
这里有个容易栽跟头的地方：默认流（也叫做 0 号流）在 CUDA 6.0 之后被分为两种行为模式，取决于你是否用 cudaStreamCreateWithFlags 加了 cudaStreamNonBlocking 标志。
| 特性 | 默认流（阻塞） | 非阻塞流 |
|---|---|---|
| 与默认流的关系 | 所有默认流操作隐式同步 | 与其他流完全隔离 |
| 对设备端全局操作的阻塞 | 阻塞全局操作（如 kernel） | 不阻塞 |
| 使用场景 | 简单脚本、调试点 | 高性能、多流重叠 |
| 创建方式 | 默认 0 流 | cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking) |
避坑：我曾经在一个多流项目里，用默认流做了数据拷贝，另一条用户流跑了内核，结果主机端一直卡在 cudaMemcpy 上。查了半天才发现，默认流的隐式同步锁住了整个设备。从那以后，我所有“并发流”都用非阻塞模式创建。
事件同步
多个流并行时，你需要一个“握手”机制。CUDA 事件就是这个握手信号。
- 记录事件：
cudaEventRecord(event, stream)。把事件插入流的末尾。 - 等待事件：
cudaStreamWaitEvent(stream, event)。当前流会在执行到这个位置时，等事件完成才继续。 - 主机同步：
cudaEventSynchronize(event)。阻塞主机直到事件背后的所有操作完成。
说实话，事件同步是处理跨流依赖的最优雅方式。你不需要全局阻塞整个设备，只需要在关键路径上插一个“哨兵”。
技巧：插事件时的开销很小，但别滥用。我见过有人在每次内核提交后都插一个 cudaEventRecord，结果事件队列爆了，反而拖慢。我的建议是：只在需要“检查点”或“拷贝依赖”的位置记录。
流优先级
现代 CUDA 架构（计算能力 3.5+）支持流的优先级。高优先级的流中的操作会优先获得计算资源。
cudaStream_t high, low;
cudaStreamCreateWithPriority(&high, cudaStreamNonBlocking, -1); // 最高优先级
cudaStreamCreateWithPriority(&low, cudaStreamNonBlocking, 0); // 普通优先级优先级数值越小越优先（嗯，别搞反了）。这个特性在混合负载中特别有用：你可以把实时推理丢进高优先级流，把日志记录或离线分析扔到低优先级流，避免推理被刷下来。
不过，优先级不是万能的。如果两个高优先级流一直在抢资源，它们之间仍然会被公平调度。说白了，优先级只是相对顺序的提示，不是硬保证。
下面这张图展示了我个人最常用的一套多流流水线结构。你可以看到，3 条流完全独立，只有事件在关键路径上同步。
异步操作的真实收益
我在生产环境测过一次：用一个 ResNet-50 推理服务，输入是 256×256 的图像。单流的极端情况——先拷贝所有图像，再跑所有推理，最后拷贝结果——延迟 72ms。而换成 4 条流 + 事件同步的流水线模式，延迟降到 31ms。说白了，异步就是让 GPU 的“搬运工”和“工人”同时干活，而不是等人来搬完才动工。
更关键的是，多流并不增加显存消耗多少（事件对象和流对象本身非常轻量）。你完全可以把所有流都创建好，在推理循环里直接提交任务，主机端只做最少的同步检查。我个人建议的规矩是：每个 GPU 至少配置 2~4 条流，对于混合负载（比如推理 + 后处理），再额外加一条高优先级流。
注意：流多了也不是越多越好。GPU 上的硬件调度器数量有限，流太多反而增加调度开销。我曾经在 V100 上试过 32 条流同时跑内核，结果性能不升反降。一般来说，流数量 = 2 × SM 数量 是比较稳妥的起点。
最后一点：事件与主机同步的取舍
如果你只是需要避免数据竞争，cudaStreamSynchronize 就够用了。但如果你关心“哪一步先完成”的时序，比如数据预处理和模型推理是并行的，事件会更细腻。我最常用的模式是：在多个流的结尾各插一个事件，然后在主机端调用 cudaEventSynchronize 等待所有事件；这样既能保持流之间的并发，又能准确知道所有结果何时就绪。
用事件同步还有一个好处：你可以用 cudaEventElapsedTime 精确统计从事件记录到完成的时间。这对于做延迟分析、性能调优非常顺手。