CUDA多卡同步、NCCL通信与性能调优 - 第2节
一、CUDA Stream 到底是什么？
说实话，我刚接触 CUDA 的时候，对 Stream 的理解就是「队列」。你往队列里丢任务，GPU 按顺序执行。但后来我发现，这个理解太浅了。
Stream 本质上是一个命令队列。GPU 会按照你提交的顺序，依次执行队列里的 kernel 启动、内存拷贝、事件记录等操作。每个 Stream 都是独立的，互不干扰。
我习惯把 Stream 想象成「流水线」。一条流水线上，工人按顺序干活。但如果你开多条流水线，它们就能同时干活——这就是多流并发的核心思想。
核心要点：Stream 是 GPU 上执行异步操作的逻辑序列。不同 Stream 之间的操作可以重叠执行，前提是硬件资源允许。
二、默认流 vs 非阻塞流
2.1 默认流（NULL Stream）
每个 CUDA 程序启动时，都会有一个默认流。你如果不显式创建 Stream，所有操作都默认提交到这个流里。
默认流有个特殊行为：它会阻塞其他流。什么意思呢？
举个例子。你在默认流里启动了一个 kernel，又在另一个流里启动了另一个 kernel。默认流里的 kernel 没跑完之前，另一个流里的 kernel 不会开始执行。这就是所谓的「隐式同步」。
我曾经踩过的坑：刚开始写多流代码时，我习惯把所有操作都丢到默认流里。结果发现 GPU 利用率上不去，性能还不如单流。后来才意识到，默认流的阻塞特性把并发全毁了。
2.2 非阻塞流（Non-blocking Stream）
要真正实现多流并发，你得用非阻塞流。创建方式很简单：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);默认情况下，用 cudaStreamCreate 创建的流就是非阻塞的。它们之间不会互相等待，除非你显式做了同步。
但要注意：非阻塞流和默认流之间，默认流仍然有阻塞行为。所以我的建议是——尽量别用默认流。所有操作都放到显式创建的流里，这样行为更可控。
个人习惯：我一般在程序初始化时就创建好所有需要的流，然后全程只用这些显式流。默认流？我当它不存在。
三、多流并发执行模式
3.1 并发能带来什么？
多流并发的核心目标是：让 GPU 的不同硬件单元同时工作。
你想想看，GPU 里有计算单元、内存拷贝引擎、还有各种专用硬件。如果只有一个流，那这些硬件可能轮流闲着。多流并发，就是让它们「各忙各的」。
举个实际场景：
- Stream A 在做矩阵乘法（计算密集型）
- Stream B 在从 CPU 拷贝数据到 GPU（内存密集型）
- Stream C 在做卷积（计算密集型，但和 A 用的硬件不同）
这三个流如果能同时跑，GPU 的利用率就上去了。
3.2 并发执行的模式
多流并发有两种常见模式：
模式一：流水线并行
把一个大任务拆成多个阶段，每个阶段放到不同的流里。比如：
// 伪代码示例
for (int i = 0; i < num_batches; i++) {
int stream_id = i % num_streams;
cudaMemcpyAsync(d_data[i], h_data[i], size, cudaMemcpyHostToDevice, streams[stream_id]);
kernel<<<grid, block, 0, streams[stream_id]>>>(d_data[i]);
cudaMemcpyAsync(h_result[i], d_result[i], size, cudaMemcpyDeviceToHost, streams[stream_id]);
}这个模式里，每个流处理一个 batch。多个流之间，内存拷贝和 kernel 执行可以重叠。
模式二：任务并行
不同流执行不同类型的任务。比如一个流专门做计算，另一个流专门做数据传输。
关键点：并发不是自动发生的。你需要确保不同流之间的操作没有数据依赖。如果有依赖，你得用事件（Event）来做同步。
3.3 并发能到多少？
这取决于你的 GPU 型号。一般来说：
| GPU 架构 | 建议流数量 | 说明 |
|---|---|---|
| Kepler (SM 3.x) | 2-4 | 并发能力有限 |
| Maxwell (SM 5.x) | 4-8 | 有所改善 |
| Pascal (SM 6.x) | 8-16 | 支持 Hyper-Q |
| Volta/Turing (SM 7.x) | 16-32 | 并发能力很强 |
| Ampere+ (SM 8.x/9.x) | 32+ | 建议 4-8 个就够 |
我个人经验：不是流越多越好。流太多，调度开销反而会拖慢性能。我一般用 4-8 个流，效果最好。
四、Stream 回调函数
4.1 回调是什么？
Stream 回调允许你在某个流里的所有操作完成后，自动执行一个 CPU 函数。说白了就是「流干完活了，通知我一声」。
用法很简单：
void CUDART_CB my_callback(cudaStream_t stream, cudaError_t status, void *user_data) {
// 这里可以做一些 CPU 端的工作
printf("Stream %p 完成了所有操作\n", stream);
}
// 在流里注册回调
cudaStreamAddCallback(stream, my_callback, NULL, 0);4.2 回调的注意事项
嗯，这里要注意几个坑：
- 回调在 CPU 上执行，不是在 GPU 上。所以别在回调里启动 kernel，那会阻塞。
- 回调函数必须是
__host__函数，不能是__device__或__global__。 - 回调的执行顺序：同一个流里，回调按注册顺序执行。不同流之间，回调的执行顺序不确定。
- 回调里不要做耗时操作，否则会阻塞后续的流操作。
我曾经犯过的错：在回调里调用了 cudaMemcpy，结果导致死锁。因为回调执行时，流还处于「忙碌」状态，你再往这个流里提交操作，就会卡住。
4.3 回调的实际用途
回调最常见的用途是：异步通知。比如：
- 数据拷贝完成后，通知 CPU 开始处理结果
- 多个流都完成后，触发一个汇总操作
- 记录性能日志，统计每个流的执行时间
我个人习惯用回调来做「轻量级的同步」。比轮询 cudaStreamSynchronize 要优雅得多。
五、知识体系总览
下面这张图，是我自己总结的多流并发知识体系。你可以把它当作本章的「地图」：
六、总结
这一章我们聊了 Stream 的核心概念、默认流与非阻塞流的区别、多流并发的两种模式，还有回调函数的使用技巧。
说实话，Stream 这个东西，光看文档是学不会的。你得动手写代码，跑 profiler，看 GPU 利用率曲线。我当年就是靠一遍遍调优，才慢慢摸清门道的。
下一章我们会深入讲「事件（Event）与同步机制」，那是多流并发的「交通指挥员」。不过那是后话了，先把这一章的内容消化掉。
最后送你一句话：多流并发不是银弹。如果你的 kernel 本身就很小，或者数据依赖很强，多流反而会引入额外开销。先分析，再优化。
想要更多资料，联系可可