并发内核执行：使用多个流实现内核并发、网格级并发与流级并发
好，咱们接着聊并发内核执行。上一章我们讲了流的基本概念，说白了流就是GPU上的一条任务流水线。但一条流水线能做的事太有限了，真正让GPU火力全开的，是让多条流水线同时跑起来——这就是多流并发。
我个人习惯把并发内核执行分成两个层次来看：网格级并发和流级并发。这两个概念容易搞混，我当年刚接触CUDA时也踩过坑，咱们今天一次说清楚。
网格级并发 vs 流级并发
先看个直观的例子。假设你有两个内核函数：kernelA和kernelB。你想让它们同时跑。
- 网格级并发：同一个流里，多个网格（内核调用）在硬件上重叠执行。嗯，这个其实很少见，因为同一个流默认是串行的。
- 流级并发：不同流里的内核，在GPU上同时执行。这才是我们真正要用的并发方式。
说白了，网格级并发是「一条流水线上多个任务挤在一起跑」，流级并发是「多条流水线同时开工」。你想想看，哪个更实用？当然是后者。
多流并发的代码骨架
来个最基础的例子。我习惯用两个流来演示：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在stream1中启动内核A
kernelA<<<grid, block, 0, stream1>>>(d_A);
// 在stream2中启动内核B
kernelB<<<grid, block, 0, stream2>>>(d_B);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);注意看，两个内核分别指定了不同的流。只要硬件资源够，它们就会在GPU上同时执行。我在项目中遇到过一种情况：两个内核一个计算密集、一个访存密集，把它们分到不同流里，吞吐量直接翻倍。
关键点：并发执行的前提是GPU有足够的计算单元和内存带宽。如果两个内核都占满SM，那它们只能排队，并发就无从谈起。
流级并发的硬件限制
这里有个坑，我曾经踩过。你以为创建了10个流，10个内核就能同时跑？太天真了。
实际上，GPU能同时执行的内核数量受硬件限制。以我的经验来看：
| GPU架构 | 最大并发内核数 | 我的备注 |
|---|---|---|
| Kepler (SM 3.x) | 16 | 老架构了，基本见不到了 |
| Maxwell (SM 5.x) | 32 | 当年我用它做过实验 |
| Pascal (SM 6.x) | 64 | 并发能力明显提升 |
| Volta/Turing (SM 7.x) | 128 | 嗯，这个数量很够用了 |
| Ampere+ (SM 8.x/9.x) | 128+ | 具体看驱动和硬件配置 |
为什么会这样？因为每个流都需要占用硬件调度资源。你创建128个流，但实际能同时跑的内核可能只有16个，剩下的都在排队。我建议：流的数量不要超过GPU的并发上限，否则白白浪费创建开销。
网格级并发的正确理解
咱们再回头说说网格级并发。这个概念其实有点误导人。在CUDA里，同一个流中的内核默认是串行执行的——前一个内核没跑完，后一个不会开始。
但有一种特殊情况：动态并行。子网格可以在父网格还没结束时启动。不过这个特性用得少，而且调试起来很痛苦。我个人建议：除非你特别清楚自己在做什么，否则别碰动态并行。
小技巧：如果你想验证两个内核是否真的在并发执行，可以用cudaEvent打时间戳。我曾经用这个方法抓到一个「假并发」的bug——两个内核其实在串行跑，只是时间太短看起来像同时执行。
实战：多流并发模式
我常用的多流并发模式有两种：
- 流水线模式：每个流处理一批数据，流之间独立。适合数据并行。
- 任务并行模式：不同流执行不同功能的内核。比如一个流做矩阵乘法，另一个流做向量归约。
来看个流水线模式的代码片段：
const int numStreams = 4;
cudaStream_t streams[numStreams];
for (int i = 0; i < numStreams; i++) {
cudaStreamCreate(&streams[i]);
}
// 每个流处理一块数据
for (int i = 0; i < numStreams; i++) {
int offset = i * chunkSize;
kernel<<<grid, block, 0, streams[i]>>>(d_data + offset);
}
// 同步所有流
for (int i = 0; i < numStreams; i++) {
cudaStreamSynchronize(streams[i]);
cudaStreamDestroy(streams[i]);
}注意看，这里每个流处理的数据是不重叠的。如果数据有依赖关系，那并发就玩不转了。我在项目中遇到过有人强行让有依赖的内核并发，结果数据全乱套了——嗯，调试了一整天。
警告：多流并发时，不要假设内核的执行顺序。不同流之间的内核执行顺序是不确定的。如果你需要顺序保证，要么用同一个流，要么用事件同步。
流级并发的性能调优
最后聊点实际的。怎么判断你的多流并发有没有生效？
我一般用nvidia-smi或者Nsight Systems看GPU利用率。如果利用率曲线是平的、接近100%，说明并发做得不错。如果曲线像锯齿一样忽高忽低，那可能有些流在空转。
还有个常见问题：内存拷贝和内核执行能不能重叠？ 答案是：可以，但需要满足条件。
- 使用固定内存（pinned memory）
- 拷贝方向与内核执行方向不同（比如拷贝到设备的同时，另一个流的内核在计算）
- 硬件支持重叠（大部分现代GPU都支持）
我曾经优化过一个视频处理管线，把数据加载、预处理、核心计算分到三个流里，配合异步拷贝，整体吞吐量提升了40%。说白了，就是让GPU的每个部件都别闲着。
好了，关于并发内核执行就聊这么多。记住一句话：流是并发的载体，但并发不是免费的。创建流有开销，同步有开销，数据依赖会限制并发。用得好，性能起飞；用不好，代码复杂度上去了，性能反而下降。
下一章咱们聊聊事件同步，那是控制流执行顺序的利器。