CUDA Stream详解

date

Jun 22, 2025

slug

cuda_stream

🥰总结

⁉️问题

如果你能快速回答出如下问题, 说明你已经掌握相关内容, 就不需要阅读本文章啦

cuda stream用在什么场景?

CUDA的异步API为什么都需要绑定一个stream?

默认流和非默认流有什么区别?

🧐内容

是什么?

我们首先看下概念: 在CUDA编程中, 流(stream)用于管理异步操作的并发执行.

这里有两个核心点: “管理异步操作”和“并发执行”

管理异步操作

我们知道GPU的一大优势就是并发度高, 任务可以异步下发, 同时执行. 但是不同的任务有可能存在依赖关系, 需要串行执行(比如先要读取数据, 等到数据读取完成再进行计算), 那么我们就需要一种资源, 能让我高效的控制任务的并行度, 这就是stream.

同一个stream里的任务是串行执行的, 不同stream里的任务是并发执行的(除了默认流, 这个比较特殊, 后面会讲到)

并发执行

首先我们需要区分两个概念: 并发和并行

概念	CUDA流中的含义	示例
并发(Concurrency)	多个操作在时间上重叠执行，但不一定同时进行（通过快速切换任务实现"同时"的假象）	CPU准备数据时GPU同时计算（通过异步操作时间重叠）
并行(Parallelism)	多个操作真正同时执行（依赖硬件能力：多SM并行/计算与复制同时进行）	多个kernel在不同流中同时运行（需Hyper-Q支持）

这个比较好理解, 我们因为要异步执行任务, 那么就需要将任务持续提交, 然后在需要用到的时候去获取结果, 不同的stream保证了我们的任务可以异步并发执行.

流首先保证并发

所有非默认流操作天然并发执行（与主机线程异步）, 即使只有1个SM的GPU，也能实现:

并行需要条件

因此，严格来说：

所有流操作本质是并发

在足够资源下升级为并行

类型

默认流

也称为空流(0或cudaStreamDefault), 当我们在调度任务时如果不指定stream, 那么就默认在该stream上执行, 比如 kernel<<<grid, block>>>() 或cudaMemcpy()

该stream有个非常重要的特性就是: 隐式同步

隐式同步

指默认流会阻塞所有其他流, 直到自身操作完成, 默认流及其他流才会继续执行. 也就是说默认流中的任何操作(如内核启动, 同步内存拷贝)都会在GPU上创建一个隐式的同步点, 具体操作如下:

暂停所有非默认流的执行

等待其他stream上已经启动的操作完成

执行默认流任务完成

恢复其他流的执行

如下图nsight所示:

非默认流

需要显示创建, 任务可以在这些流中并行执行, 主要用于异步任务的执行

常见用法

流同步

显式同步, 通常用于当前流的同步

事件(Event), 通常用于不同流之间的同步

优先级

高优先级流可优先占用 GPU 资源：

最佳实践

避免隐式同步

对于一些简单的程序, 默认流的隐式同步能够简化代码, 不需要用户手动控制流. 但是对于复杂程序, 我们可能会使用多个流, 默认流的隐式同步操作可能会导致不必要的性能问题(比如我希望任务能完全并行执行, 中间没有同步操作), 这时我们就不需要这个隐式同步了. CUDA也提供了解决办法.

通过以下接口创建流:

cudaStreamNonBlocking 是创建流时的标志参数，其核心作用是解除新建流与默认流的隐式同步关系。这解决了传统非默认流的一个关键性能瓶颈。

两种参数对比:

Flag类型	同步行为	性能影响	适用场景
cudaStreamDefault	隐式与默认流同步	并行性受限	兼容旧代码
cudaStreamNonBlocking	完全独立于默认流	最大化并行性能	高性能计算, 流水线优化

计算与数据拷贝重叠

多设备协同

结合 cudaSetDevice() 为每个设备分配独立流：