4. 流与并发执行：CUDA流(Stream)概念、默认流与非默认流、流同步机制、多流并发执行模式
好，咱们今天聊一个实战中特别关键的话题——CUDA流。说白了，流就是GPU上的一串操作队列。你往队列里丢任务，GPU就按顺序执行。但如果你只开一个队列，那GPU的很多计算单元就闲着，效率上不去。
我个人习惯把流想象成「流水线」。一条流水线干活，另一条也能同时干活。这就是并发的核心。
4.1 什么是CUDA流？
CUDA流，本质上是一个异步操作的序列。同一个流里的操作，严格按照顺序执行。不同流里的操作，则可以乱序、甚至并发执行。
举个例子：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在stream1中启动kernel
kernel1<<<grid, block, 0, stream1>>>(d_data1);
// 在stream2中启动kernel
kernel2<<<grid, block, 0, stream2>>>(d_data2);
这里kernel1和kernel2如果硬件资源够，就会同时跑。我在项目中遇到过，有些新手把所有操作都丢到默认流里，结果GPU利用率只有30%。一拆成多流，直接飙到80%以上。
4.2 默认流 vs 非默认流
每个CUDA程序启动时，都会自动创建一个默认流。你没显式指定流的时候，所有操作都默认进这个流。
这里有个坑——默认流在某些设备上是阻塞的。什么意思？就是默认流里的操作，会阻塞其他流里的操作。嗯，这其实是个历史遗留问题。
| 特性 | 默认流（NULL流） | 非默认流 |
|---|---|---|
| 创建方式 | 自动存在 | cudaStreamCreate() |
| 阻塞行为 | 可能阻塞其他流 | 不阻塞其他流 |
| 并发能力 | 弱 | 强 |
| 使用场景 | 简单测试 | 生产环境 |
4.3 流同步机制
多流并发虽好，但你不能让它们乱跑。比如stream1算完了结果，stream2要用这个结果。这时候就需要同步。
常用的同步方式有几种：
- cudaDeviceSynchronize()：阻塞CPU，直到所有流都完成。简单粗暴，但效率低。
- cudaStreamSynchronize(stream)：只等某一个流完成。
- cudaEventSynchronize(event)：用事件做更细粒度的同步。
我给你们看个实际例子：
cudaEvent_t event;
cudaEventCreate(&event);
// stream1做完后，记录事件
kernel1<<<grid, block, 0, stream1>>>(d_data);
cudaEventRecord(event, stream1);
// stream2等event完成后再执行
cudaStreamWaitEvent(stream2, event, 0);
kernel2<<<grid, block, 0, stream2>>>(d_data);
我曾经在一个视频处理项目里，用事件同步把流水线延迟从5ms降到了1.2ms。你想想看，这差距有多大。
4.4 多流并发执行模式
多流并发，不是简单地把任务扔到不同流里就完事了。你得考虑硬件资源。
GPU内部有多个引擎：
- 计算引擎（SM）
- 拷贝引擎（DMA）
- 显存控制器
这些引擎是可以并行工作的。所以多流并发的最佳实践是：
- 计算与拷贝重叠：一个流在做数据传输，另一个流在做计算。
- 多计算流并行：多个小kernel同时跑，填满SM。
- 流水线模式：把一个大任务拆成多个阶段，每个阶段一个流。
我给你们画个典型的多流并发模式：
// 流0：负责数据传输
cudaMemcpyAsync(d_in1, h_in1, size, cudaMemcpyHostToDevice, stream0);
// 流1：负责计算
kernel1<<<grid, block, 0, stream1>>>(d_in1, d_out1);
// 流2：负责结果回传
cudaMemcpyAsync(h_out1, d_out1, size, cudaMemcpyDeviceToHost, stream2);
这里有个细节——同一个流里的操作是串行的。所以你要把不依赖的操作放到不同流里。依赖的操作，用事件同步。
4.5 避坑指南
讲几个我踩过的坑：
- 流数量不是越多越好。GPU的硬件队列有限，一般4-16个流就够了。我曾经试过开32个流，结果性能反而下降了。
- 小心隐式同步。比如cudaMalloc、cudaMemcpy（不带Async的版本）都会隐式同步所有流。我建议统一用cudaMemcpyAsync。
- 事件要重用。别每次同步都创建新事件，用cudaEventRecord反复记录就行。
嗯，说到隐式同步，我记得有一次排查性能问题，找了半天发现是cudaMalloc在作怪。每次分配显存都让所有流停下来等。后来改成预分配，性能直接翻倍。
4.6 总结
流与并发执行，说白了就是让GPU忙起来。别让它闲着。默认流是新手村，多流才是高手区。
我建议你从两个流开始练手：一个做数据传输，一个做计算。等跑通了，再慢慢加流。记住，并发不是目的，提高吞吐才是。
下一章咱们聊内存优化，那又是另一个大坑。到时候我给你们讲讲怎么用统一内存和零拷贝，让数据传输不再是瓶颈。