做AI推理的时候，如果你用的是NVIDIA显卡加上TensorRT，那一定遇到过这样的场景：程序里同时来了好几个用户的请求，每个请求都要跑一遍模型。最笨的办法就是排队，一个一个来，简单是简单，但GPU的利用率可能连一半都不到。想要让这些请求同时跑起来，就得开多条“流水线”，也就是多条CUDA流。今天这篇文章，就专门聊聊多流执行时的时序问题，特别是流与流之间的同步，以及CUDA事件到底怎么用才不容易踩坑。

一、应用场景：什么时候需要多流执行？

先别急着看代码，我们先想清楚一个问题：什么情况下你才需要用到多流？

最常见的场景就是在线服务。比如一个图像识别的API，同时收到了十几个用户的图片，每张图片都要经过同一个模型得到结果。如果用单流，那就是一个人把十几张图挨个处理完才能给下一批人办理，后面的人等到花儿都谢了。而多流就像食堂里多开了几个打饭窗口，每个人排不同的队列，整体效率自然高很多。

还有一种场景是“一个请求里面包含了多个不同的模型”。典型的例子是视频分析：先要用一个小模型检测出画面里的物体，再把物体区域交给一个更精细的分类模型。这两个模型愿意的话，可以分别放在两条流里，让GPU同时干活，节省总的响应时间。

另外，预处理和后处理也能和推理重叠起来。比如从硬盘读图片、缩放到模型需要的大小，这些操作如果和推理放在同一条流里，那GPU得等着CPU先干完。如果单独开一条流做预处理，主流的推理就可以同时进行。所有这些需求，最后都会指向同一个工具——CUDA流。

二、先搞懂底层的两条腿：流和CUDA事件

要想会用多流，脑子里得有一个具体画面。我们把GPU想象成一个大型厨房，里面有很多厨师（计算单元）。每条流就是一条装着各种食材和任务单的传送带。传送带上的任务必须按顺序执行，这是流的特点。不过，不同的传送带之间是互不影响的，可以同时往前走。所以多流的本质就是：让厨房里不同的厨师，同时处理不同传送带上的菜。

那CUDA事件又是什么呢？它更像一个“对讲机”或者“哨子”。当一条流把某个关键步骤做完时，它可以通过事件把这个消息广播出去。另一条流如果想确保这一步骤已经完工，就可以在合适的地方停下来，等着听这个哨子声。有了事件，流和流之间才能协调节奏，避免互相踩脚。

这里要特别注意：流本身是异步的，你发出一个任务后，代码不会等你执行完再走。如果你想拿到结果，必须主动去同步。CUDA事件就是用来做这种“主动同步”的工具之一。

三、多流执行中的三个大坑

生活中有句话叫“好心办坏事”，多流用不好，也会让自己掉进坑里。下面这几种情况，我身边的朋友基本都碰到过。

3.1 坑一：所有流操作共用了同一块显存

这是最隐蔽也最容易犯的错。你开了两条流，为了省事，给它们分配了同一个输入缓冲区。结果第一个请求还没读完，第二个请求就冲进来把数据改了。等你把推理结果拿出来一看，全乱了。

要解决这个问题，最直接的办法就是每条流使用自己独立的输入输出内存，也就是给每个请求都单独准备一套“碗筷”。千万别抠门，显存分配那点开销，比起数据错乱带来的调试痛苦，真的不算什么。

3.2 坑二：同步手段用错，该等的没等，不该等的等了

有一些开发者，包括我以前也这么干过：为了拿结果，直接调用cudaDeviceSynchronize()。这个函数会死死地等住整个设备，直到所有流的活都干完。这样一来，本来想并行加速的，结果全被这个笨重的同步给拽回单流了。这就是“不该等的等了”。

反过来，还有“该等的没等”。比如你让一条流执行推理，然后连一次同步都不做，立刻在CPU侧去读输出缓冲区的数据。这时候推理很可能还没跑完，你读到的只是一堆乱码或者上一次的旧数据。正确的做法是，在读取数据之前，先让对应的流完成。你可以用cudaStreamSynchronize()，或者用后面要讲的事件。

3.3 坑三：CUDA事件被用成了“死亡锁”

事件虽然好用，但用错时机就会让程序卡住。最典型的错误是“自己等自己”：在一条流里记录了一个事件，然后马上让同一条流去等这个事件。这就好比一个人站在传送带前面，等传送带把自己送过去，可传送带又因为他站着不动而卡住了，两边互相瞪着，最后死锁。

还有一种情况：事件在流A上记录，但你在流B上等它，同时流A后续还要等流B的另外一个事件。这种互相等待一旦形成环，整个推理管线就瘫了。所以记住，事件等待的方向要清楚，别在你的流之间搞出循环依赖。

四、一个安全可靠的多流推理示例

下面我用一个完整的Python示例，演示怎么用两条流配合CUDA事件，完成两路推理。这里假设第二个推理必须等第一个推理完成才能开始（比如是级联模型），但两个流的输入拷贝可以并行。

# 技术栈：Python（配 TensorRT 和 PyCUDA）
import tensorrt as trt
import pycuda.driver as cuda
import pycuda.autoinit
import numpy as np

# 假设你已经有了一个 TensorRT engine 文件
TRT_LOGGER = trt.Logger(trt.Logger.WARNING)
with open("model.engine", "rb") as f:
    runtime = trt.Runtime(TRT_LOGGER)
    engine = runtime.deserialize_cuda_engine(f.read())

# 创建两个 CUDA 流：主流和辅助流
stream_main = cuda.Stream()
stream_helper = cuda.Stream()

# 创建两个执行上下文，每个流配一个，防止状态串台
ctx1 = engine.create_execution_context()
ctx2 = engine.create_execution_context()

# 获取输入输出张量的形状（假设张量名叫 input 和 output）
input_shape = engine.get_tensor_shape("input")
output_shape = engine.get_tensor_shape("output")

# 为两条流分别准备主机侧的内存（用分页锁定内存，方便异步拷贝）
h_in1 = cuda.pagelocked_empty(trt.volume(input_shape), dtype=np.float32)
h_out1 = cuda.pagelocked_empty(trt.volume(output_shape), dtype=np.float32)
h_in2 = cuda.pagelocked_empty(trt.volume(input_shape), dtype=np.float32)
h_out2 = cuda.pagelocked_empty(trt.volume(output_shape), dtype=np.float32)

# 为两条流分别分配设备侧显存
d_in1 = cuda.mem_alloc(h_in1.nbytes)
d_out1 = cuda.mem_alloc(h_out1.nbytes)
d_in2 = cuda.mem_alloc(h_in2.nbytes)
d_out2 = cuda.mem_alloc(h_out2.nbytes)

# 创建两个事件，用于流间同步
event_first_done = cuda.Event()
event_second_done = cuda.Event()

# ---------- 第一路推理（主流） ----------
# 把输入数据从主机内存异步拷贝到显存
cuda.memcpy_htod_async(d_in1, h_in1, stream=stream_main)
# 在主流上启动推理
ctx1.execute_async_v2(
    bindings=[int(d_in1), int(d_out1)],
    stream_handle=stream_main.handle
)
# 记录事件：主流已经完成了第一路推理
event_first_done.record(stream_main)

# ---------- 第二路推理（辅助流） ----------
# 辅助流先把第二路输入拷贝好，这个动作不需要等主流
cuda.memcpy_htod_async(d_in2, h_in2, stream=stream_helper)
# 关键：辅助流等待第一路推理完成的事件
stream_helper.wait_for_event(event_first_done)
# 等到了事件后，辅助流再执行第二路推理
ctx2.execute_async_v2(
    bindings=[int(d_in2), int(d_out2)],
    stream_handle=stream_helper.handle
)
# 记录辅助流完成事件
event_second_done.record(stream_helper)

# ---------- 拷回结果 ----------
# 注意：拷贝结果时也要放在对应的流里，不能直接在CPU侧乱读
cuda.memcpy_dtoh_async(h_out1, d_out1, stream=stream_main)
stream_helper.wait_for_event(event_second_done)
cuda.memcpy_dtoh_async(h_out2, d_out2, stream=stream_helper)

# 最后，让两条流都彻底干完，保证拷贝也完成
stream_main.synchronize()
stream_helper.synchronize()

# 打印一下前5个数，证明数据真的回来了
print("第一路输出前5个值:", h_out1[:5])
print("第二路输出前5个值:", h_out2[:5])

这段代码里，最值得记在小本本上的两句话是：event_first_done.record(stream_main) 和 stream_helper.wait_for_event(event_first_done)。前者表示“我在主流这一步做完时吹个哨子”，后者表示“辅助流先停下，听到哨子再走”。这个模式非常常见，多模型级联、多请求批次对齐，都能套用。

有人可能会问：既然第二路要等第一路，那这两条流还有什么并行可言？注意看，第二路在等待事件之前，先把输入拷贝发出去了。这块拷贝操作和第一路的推理是重叠的，所以仍然比完全串行快。如果你有两个完全独立的请求，甚至可以把中间那次等待删掉，两条流彻底平行跑，这样并行的收益更大。

五、技术优缺点和注意事项

5.1 多流执行的优势

最大的优点当然是提升GPU利用率。模型推理的时候，往往不是所有计算单元都被塞满，多流可以让空闲的单元去处理别的请求，从而提升整体的吞吐量。另外，多流还能把数据拷贝和计算重叠起来，减少总耗时的感觉。

5.2 多流执行带来的代价

代价也有，而且不小。首先，你得为每条流准备独立的内存和上下文，这让内存开销变大。其次，代码复杂度上了一个台阶，什么时候该等，什么时候不该等，心里得像明镜一样。一旦同步失误，出现的都是很难复现的随机错误。

5.3 注意事项清单

  • 每条流尽量使用独立的执行上下文，不要共用一个。
  • 输入输出显存要按流隔离，别贪图省事共用同一块。
  • 等待事件之前，确认事件已经在另一条流上记录过了。
  • 绝对不要在同一个流里记录事件后立刻等同一个事件，会死锁。
  • 读取结果前，记得用流同步或者事件同步，确保数据已经就绪。
  • 不要滥用cudaDeviceSynchronize，它会打断所有流的并行工作。
  • 如果只是一个模型连着跑很多次，有时候单流加批处理可能比多流更简单，性能也不差。多流解决的是“异构任务”之间的并发问题。

六、总结

多流执行是TensorRT推理管线里的一把双刃剑。用好了，可以让GPU忙得有条不紊；用不好，轻则性能退回原点，重则程序卡死或者数据错乱。核心要点其实就两条：一是每条流的资源要隔离开，二是用CUDA事件来精确控制流之间的依赖关系。记住“资源隔离、事件同步”这八个字，下次再碰到多流问题，你心里就有底了。写代码的时候别嫌麻烦，把同步逻辑画在纸上，搞清楚了再动手，踩坑的概率会小很多。