1. CUDA并发执行模型总览
在CUDA编程的早期阶段，大多数开发者习惯将主机端代码与设备端核函数交替排列：拷贝数据、启动核函数、等待完成、再拷贝回来。这种“串行”模式在单任务场景下逻辑清晰，但当程序中出现多个相互独立的计算任务时，GPU的利用率会急剧下降。原因在于，一个核函数执行期间，复制引擎和计算单元往往处于闲置状态；而一次内存拷贝进行时，SM阵列又无事可做。CUDA并发执行模型的核心目标，就是让这些不同性质的硬件资源尽可能同时忙碌起来。
硬件队列与软件流的关系
要理解多流并发，必须先看清硬件层面的执行队列。GPU内部并非只有一个队列，而是由多个硬件工作队列组成，分别对应不同类型的操作：计算工作队列负责核函数调度，复制队列负责主机与设备之间的数据传输，还有专门处理设备内部拷贝的引擎。这些队列由GPU的硬件调度器管理，彼此独立推进。
软件层面的“流”（Stream）本质上是对这些硬件队列的抽象。一个流是一串按顺序执行的操作序列，流内操作严格保序，流间操作则可以并发。当我们创建多个流并将不同任务分配到不同流中时，驱动程序会尝试将这些操作提交到不同的硬件队列上，从而让计算与传输重叠、让多个核函数并行执行。
异步执行的价值与代价
异步执行带来的收益是直接的。假设有三个独立任务，每个任务需要先做一次主机到设备的拷贝（约2ms），再执行一个核函数（约3ms）。在串行模式下，总时间约为 3×(2+3)=15ms。如果使用三个流并发，拷贝和计算可以重叠，理想情况下总时间约为 max(3×2, 3×3)=9ms，加速比接近1.67倍。这个估算假设拷贝引擎和计算引擎完全独立且带宽充足，实际中还会受到PCIe带宽、SM占用率等因素影响，但重叠带来的收益是真实存在的。
然而，异步执行并非没有代价。首先，代码复杂度上升：需要显式创建流、管理流间依赖、处理错误传播。其次，资源竞争可能引发新的瓶颈：多个流同时进行拷贝时，PCIe带宽会被瓜分；多个核函数同时驻留时，SM的寄存器和共享内存可能成为限制。此外，流间同步如果处理不当，会出现数据竞争或死锁。因此，多流并发的正确使用，建立在对硬件队列容量、引擎数量和资源限制的准确理解之上。
从串行到并发的演进路径
实际工程中，从串行到多流并发通常经历三个阶段。第一阶段是默认流串行：所有操作都在默认流中顺序执行，逻辑最简单，但GPU利用率最低。第二阶段是手动多流：开发者创建多个非默认流，将独立任务分配到不同流中，利用流间并发实现重叠。第三阶段是流水线化：将任务进一步拆分为更细粒度的阶段，配合事件（Event）进行流间同步，形成稳定的流水线，使拷贝、计算、回传持续重叠。
以向量加法为例，假设有4个独立的数据块需要处理。串行方式下，依次执行4次“拷贝→计算→回传”，总时间约为4×(t_copy + t_kernel + t_copy)。多流方式下，创建4个流，每个流负责一个数据块的完整流程，由于不同流的拷贝和计算可以重叠，总时间可缩短至约 t_copy + 4×t_kernel + t_copy（假设拷贝引擎和计算引擎完全并行）。如果进一步使用流水线，将每个数据块再拆分为多个子块，重叠效果会更加充分。
// 多流并发示例：4个流处理4个独立数据块
const int N = 4;
cudaStream_t streams[N];
for (int i = 0; i < N; i++) {
cudaStreamCreate(&streams[i]);
}
for (int i = 0; i < N; i++) {
cudaMemcpyAsync(d_in[i], h_in[i], size, cudaMemcpyHostToDevice, streams[i]);
kernel<<<grid, block, 0, streams[i]>>>(d_in[i], d_out[i]);
cudaMemcpyAsync(h_out[i], d_out[i], size, cudaMemcpyDeviceToHost, streams[i]);
}
for (int i = 0; i < N; i++) {
cudaStreamSynchronize(streams[i]);
cudaStreamDestroy(streams[i]);
}这段代码中，每个流内的操作保持顺序，但不同流之间的拷贝和计算可以并发。需要注意的是，cudaMemcpyAsync 必须配合页锁定内存（pinned memory）才能实现真正的异步传输；如果使用普通可分页内存，驱动程序可能会退化为同步拷贝，重叠效果大打折扣。
并发执行的边界条件
多流并发并非无条件生效。以下条件必须同时满足，重叠才可能发生：第一，操作之间没有数据依赖，或者依赖已通过事件正确同步；第二，硬件具备足够的复制引擎和计算队列；第三，内存传输使用页锁定内存；第四，核函数的资源占用允许与其他核函数同时驻留。如果核函数占用了全部SM资源，后续核函数只能等待，流间并发退化为流内串行。
因此，在设计和优化多流程序时，建议先通过CUDA事件或Nsight Systems等工具测量实际重叠情况，确认瓶颈究竟在拷贝、计算还是同步上，再针对性地调整流数量、任务粒度和资源分配。盲目增加流数量不仅不会提升性能，反而可能因调度开销和资源竞争导致性能下降。