多流并发原理：GPU硬件资源、并发执行条件与流调度机制

要真正理解多流（Multi-Stream）并发，不能只停留在“创建多个流就能加速”的直觉层面。GPU 的并发能力由硬件资源决定，而流只是软件层面的“任务编排容器”。本节从硬件资源模型出发，厘清并发执行的条件，再深入流调度器的工作机制，帮助你建立从硬件到软件的正确映射。

一、GPU 硬件资源模型：并发的基础

现代 NVIDIA GPU（如 Ampere、Hopper、Ada 架构）内部包含大量可并行执行的计算单元，其核心资源可归纳为三类：

SM（Streaming Multiprocessor）：GPU 的基本计算单元，每个 SM 包含若干 CUDA Core、Tensor Core、SFU（特殊函数单元）和共享内存。一个 GPU 通常有几十到上百个 SM。
执行引擎（Execution Engine）：负责执行特定类型的操作，包括 计算引擎（执行 kernel）、拷贝引擎（DMA，负责 H2D/D2H/设备内拷贝）、以及可选的 NVENC/NVDEC 等专用引擎。
内存层次：全局显存（VRAM）、L2 缓存、SM 内的共享内存和寄存器文件。内存带宽和缓存容量直接影响并发时的数据吞吐。

并发执行的本质是：不同引擎或不同 SM 上的任务在同一时刻独立运行。例如，拷贝引擎在搬运数据时，计算引擎可以同时执行 kernel；两个 kernel 如果使用不同的 SM 子集，也可以真正并行。

关键概念：GPU 的并发粒度是“引擎级”和“SM 级”。流并发主要利用的是引擎级并行（拷贝与计算重叠）以及 SM 间的空间并行（多个 kernel 分占不同 SM）。

二、并发执行的条件：不是“创建流”就能并行

多流并发能否真正发生，取决于以下四个硬性条件，缺一不可：

条件	说明	反例
1. 资源充足	SM 数量、寄存器文件、共享内存必须能同时容纳多个 kernel	kernel 占用 100% 寄存器，第二个 kernel 必须等待
2. 引擎独立	操作必须使用不同的执行引擎（如拷贝 vs 计算）	两个 kernel 都占满所有 SM，无法并行
3. 无数据依赖	流之间的操作不能有先后依赖关系	Stream B 需要 Stream A 的输出数据
4. 显存带宽有余量	并发访问未超过显存带宽上限	两个流同时大量读写显存，带宽饱和

常见误区：很多人以为“创建多个流 = 自动并行”。实际上，如果 kernel 占满所有 SM，或者流之间存在隐式同步（如默认流与阻塞流），并发根本不会发生。并发是硬件资源调度的结果，不是软件 API 调用的必然结果。

三、流调度机制：从提交到执行

CUDA 运行时维护一个流调度器，其工作流程如下：

任务提交：主机端通过 cudaStreamCreate 创建流，并通过 cudaMemcpyAsync、kernel<<>> 等 API 将操作提交到指定流。
依赖分析：调度器检查每个操作的前置依赖（事件、流内顺序、跨流事件等待）。只有依赖满足的操作才进入“就绪队列”。
资源分配：就绪操作按优先级（高优先级流优先）和资源可用性分配到具体的执行引擎或 SM 子集。
执行与切换：引擎开始执行；当操作阻塞（如等待显存分配）时，调度器切换到其他就绪操作，保持引擎忙碌。
完成通知：操作完成后，通过事件或回调通知依赖它的后续操作。

调度器的工作核心是保持引擎尽可能忙碌。如果某个流因数据未就绪而阻塞，调度器会立即将引擎分配给其他流的就绪任务，这就是并发提升吞吐量的本质。

四、案例：典型多流并发场景

以 TensorRT 推理中的经典“流水线并行”为例，假设有 3 个输入 batch 需要处理：

// 伪代码：三流流水线
cudaStream_t s1, s2, s3;
cudaStreamCreate(&s1); cudaStreamCreate(&s2); cudaStreamCreate(&s3);

// 每个流独立执行：H2D → 推理 → D2H
for (int i = 0; i < 3; i++) {
    cudaMemcpyAsync(d_input[i], h_input[i], size, cudaMemcpyHostToDevice, streams[i]);
    myTRTInference(d_input[i], d_output[i], streams[i]);  // TensorRT 执行
    cudaMemcpyAsync(h_output[i], d_output[i], size, cudaMemcpyDeviceToHost, streams[i]);
}

理想情况下，时间线如下：

t0~t1：Stream 1 执行 H2D 拷贝（拷贝引擎忙）
t1~t2：Stream 1 推理（计算引擎忙），Stream 2 执行 H2D（拷贝引擎忙）
t2~t3：Stream 1 D2H（拷贝引擎），Stream 2 推理（计算引擎），Stream 3 H2D（拷贝引擎）

这样，拷贝引擎和计算引擎始终处于工作状态，总耗时从串行的 9 个时间段缩短为 5 个时间段左右，吞吐量提升接近 2 倍。

五、误区与复盘

复盘要点：

误区 1：流越多越好。实际上，流数量超过引擎数量后，收益递减，且增加调度开销。通常 2~4 个流即可覆盖大多数场景。
误区 2：忽略默认流的隐式同步。默认流（stream 0）与所有阻塞流有隐式同步，混用会破坏并发。建议使用非阻塞流（cudaStreamCreateWithFlags 指定 cudaStreamNonBlocking）。
误区 3：不检查资源占用。如果 kernel 的寄存器或共享内存使用量过大，SM 无法同时容纳多个 block，并发自然失效。可用 cudaOccupancyMaxPotentialBlockSize 检查占用率。
误区 4：忽略事件同步。跨流依赖必须用 cudaEventRecord 和 cudaStreamWaitEvent 显式表达，否则可能产生数据竞争。

总结：多流并发的核心是理解 GPU 的硬件资源模型——引擎独立、SM 分区、内存带宽有限。流调度器负责在满足依赖的前提下最大化资源利用率。实际开发中，应通过 profiling（如 Nsight Systems）观察引擎活动时间线，确认并发是否真正发生，再针对性地调整流数量和资源占用。