2、CUDA流（Stream）基础：流的创建、销毁与同步

好，咱们今天聊聊CUDA流。说白了，流就是GPU上的一串操作队列。你往里面丢kernel、丢内存拷贝，它按顺序执行。但多个流之间，是可以并行的。

我刚开始接触CUDA时，总觉得一个程序一个流就够了。后来做视频处理，发现GPU利用率老上不去，一查——哦，原来所有操作都在一个流里排队呢。那感觉，就像单车道堵车，明明旁边还有条空车道。

2.1 什么是CUDA流？

CUDA流，本质上是一个异步操作序列。同一个流里的操作，严格按照FIFO顺序执行。不同流之间，没有顺序保证——它们可以交错、可以并行。

你想想看，这像什么？像多个流水线。每个流水线有自己的工人（SM），有自己的任务队列。只要资源够，两条线就能同时干活。

核心要点：

  • 流内操作：串行，按提交顺序执行
  • 流间操作：无顺序依赖，可并行执行
  • 默认流（NULL stream）是同步的，会阻塞其他流

这里有个坑，我当年踩过——默认流的行为很特殊。它会跟所有非默认流同步。什么意思？就是你在默认流里提交一个操作，它会等所有其他流里的操作完成，然后才执行。反过来也一样。这其实破坏了并行的初衷。

⚠️ 避坑指南：我曾经在项目里混用默认流和自定义流，结果性能反而下降了。排查了半天，才发现是默认流的隐式同步在作怪。后来我养成了习惯：要么全用默认流，要么全用自定义流，绝不混用。

2.2 流的创建与销毁

创建流很简单，就一个API调用：

cudaStream_t stream;
cudaError_t err = cudaStreamCreate(&stream);

嗯，这里要注意。每次创建流，GPU都会分配一些资源。我见过有人在一个循环里创建了上千个流，结果显存爆了。流的数量不是越多越好，一般4-8个就够用了。

销毁流更简单：

cudaError_t err = cudaStreamDestroy(stream);

但有个细节——销毁流时，如果流里还有没执行完的操作，它会等所有操作完成再销毁。这是阻塞的。所以，我建议在销毁前先同步一下，确保流是空闲的。

💡 个人习惯：我一般这样管理流的生命周期：

  1. 程序启动时，一次性创建所有需要的流
  2. 运行期间，复用这些流，不创建也不销毁
  3. 程序退出前，统一同步并销毁

这样避免了频繁创建销毁的开销，也减少了资源泄漏的风险。

2.3 流的同步机制

同步，说白了就是「等」。你让GPU干活，但CPU可能不等它干完就继续往下走了。这时候就需要同步。

CUDA提供了几种同步方式：

同步方式 作用范围 是否阻塞CPU
cudaDeviceSynchronize() 整个设备 是
cudaStreamSynchronize(stream) 指定流 是
cudaEventSynchronize(event) 指定事件 是
cudaStreamWaitEvent(stream, event) 流等待事件 否

你可能会问：「这么多同步方式，我该用哪个？」

我的经验是：能用细粒度的，就别用粗粒度的。cudaDeviceSynchronize() 是核弹级别的，它会等所有流、所有操作都完成。这在调试时有用，但生产代码里尽量少用。

cudaStreamSynchronize() 更精准一些，只等一个流。但如果你有多个流在并行，等一个流可能意味着其他流也在等——因为资源被占用了。

我个人最常用的是事件同步。事件就像个标记，你把它插到流里，然后在另一个流里等它。这样能做到流间的精确同步，而不阻塞CPU。

2.4 实战：多流并行示例

光说不练假把式。咱们看个实际例子：

// 创建两个流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 在stream1中执行kernel A
kernel_A<<<grid, block, 0, stream1>>>(d_data1);

// 在stream2中执行kernel B
kernel_B<<<grid, block, 0, stream2>>>(d_data2);

// 等两个流都完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

// 清理
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

这段代码里，kernel A和kernel B是并行执行的。前提是GPU有足够的资源（SM、显存带宽等）。如果资源不够，它们会排队，但顺序是不确定的。

关键点：并行不是免费的。每个流都需要独立的资源。如果你的kernel已经把GPU占满了，开再多的流也没用。我做过测试，在满负载的情况下，4个流和8个流的性能几乎一样。

2.5 流的优先级

CUDA 8.0之后，流可以设置优先级。高优先级的流会优先获得资源。

int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);

cudaStream_t high_prio_stream;
cudaStreamCreateWithPriority(&high_prio_stream, 
                             cudaStreamNonBlocking, 
                             priority_high);

这个功能我用的不多，因为大多数场景下，流的优先级差异不大。但有一种情况很管用——你有一个延迟敏感的任务（比如用户交互响应），和一批后台计算任务。这时候给交互任务设高优先级，能明显提升用户体验。

⚠️ 注意：优先级不是绝对的。它只是提示调度器，不是强制命令。而且，不同GPU对优先级的支持程度不同。老一点的卡（计算能力3.x以下）可能根本不支持。

2.6 流的常见误区

我总结几个新手容易犯的错：

  • 误区一：以为流越多性能越好。实际上，流的数量超过SM数量后，收益递减。
  • 误区二：忘记检查返回值。cudaStreamCreate可能失败，尤其是显存不足时。
  • 误区三：在流之间共享数据时不加同步。不同流之间没有顺序保证，直接读写同一块内存会导致数据竞争。
  • 误区四：混用默认流和自定义流。默认流的隐式同步会破坏并行性。

我曾经在优化一个图像处理管线时，踩了误区三的坑。两个流分别处理图像的不同部分，结果一个流写数据时，另一个流正在读——画面直接花了。后来加了事件同步才解决。

2.7 知识体系总览

下面这张图，是我梳理的CUDA流知识体系。你可以把它当个地图，学完这章后对照着回顾一下：

CUDA流知识体系 核心概念 流内串行，流间并行 异步操作队列 生命周期管理 cudaStreamCreate cudaStreamDestroy 复用优于频繁创建 同步机制 设备级同步 流级同步 事件同步（推荐） 高级特性 流优先级 非阻塞流标志 常见误区 流越多性能越好 ✗ 忽略返回值检查 ✗ 混用默认流 ✗ 最佳实践 4-8个流足够 事件同步优于全局同步 统一使用自定义流

嗯，到这，CUDA流的基础知识就讲完了。创建、销毁、同步，这三个操作你得多练练。尤其是同步，它是多流编程的命门——同步少了出bug，同步多了没性能。

我个人建议，刚开始接触多流编程时，先用事件同步把逻辑理清楚。等熟练了，再考虑用cudaStreamWaitEvent做更精细的控制。别一上来就追求极致性能，先把正确性保证好。

💡 一个小技巧：调试多流程序时，可以用cuda-gdb的info streams命令查看所有流的状态。这比printf好用多了。

公众号:蓝海资料掘金营，微信deep3321

← 上一章 📖 返回目录 下一章 →