4、流调度机制：GPU硬件如何并行执行多个流
好，咱们接着聊。前面我们讲了流的概念，也讲了怎么创建和使用流。但有个问题一直悬着——你往两个流里扔了任务，GPU到底是怎么并行执行的？是真正同时跑，还是只是看起来像并行？
这就要深入到硬件层面了。说白了，流调度机制就是GPU内部的交通指挥系统。我当年刚接触CUDA时，以为只要创建了多个流，性能就能翻倍。结果一测，发现两个流跑出来的时间跟一个流差不多。当时我就懵了——这硬件到底在搞什么鬼？
4.1 GPU的硬件队列结构
先看看GPU内部的结构。每个GPU都有多个硬件队列，这些队列才是真正干活的东西。
关键概念：CUDA流是软件抽象，硬件队列是物理实体。一个流可以映射到多个硬件队列，但一个硬件队列同一时刻只能处理一个流里的任务。
我习惯把GPU想象成一个大型工厂。工厂里有好几条生产线（硬件队列），每条生产线都有专门的工人（计算单元）。你创建的流，就像是给这些生产线下的订单。
具体来说，NVIDIA GPU内部有这些队列：
- 计算队列：处理kernel执行
- 拷贝队列：处理数据传输
- 同步队列：处理事件和依赖关系
嗯，这里要注意——不同代的GPU，硬件队列数量不一样。比如Kepler架构有32个硬件队列，而Pascal之后变成了更多。你想想看，如果你创建了64个流，但硬件只有32个队列，那剩下的流就得排队等着。
4.2 流调度的核心机制：超线程与并发
GPU怎么做到多个流并行？核心机制是超线程调度（Hyper-Q）。
Hyper-Q是NVIDIA在Kepler架构引入的技术。它允许CPU在同一个GPU上同时提交多个独立的任务流。以前的老架构（Fermi）只有一个硬件工作队列，所有流都得排队。Hyper-Q把这个队列扩展到了32个。
我举个例子你就明白了：
// 没有Hyper-Q的情况（Fermi架构）
cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);
kernel1<<<grid, block, 0, s1>>>(...);
kernel2<<<grid, block, 0, s2>>>(...);
// 实际上kernel2必须等kernel1执行完才能开始为什么会这样？因为只有一个硬件队列，所有任务都塞进去了。就像只有一个收银台的超市，不管你有多少购物车，都得一个一个结账。
但有了Hyper-Q之后：
// 有Hyper-Q的情况（Kepler及之后）
cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);
kernel1<<<grid, block, 0, s1>>>(...);
kernel2<<<grid, block, 0, s2>>>(...);
// 两个kernel可以同时执行，只要硬件资源够用我在项目中遇到过一个问题：明明用了多个流，但性能提升不明显。后来发现是kernel太小了，每个kernel的执行时间比调度开销还短。这种情况下，并行反而更慢。
4.3 硬件调度器的决策逻辑
GPU的硬件调度器（Warp Scheduler）是怎么决定先执行哪个任务的？
说白了，它看的是资源可用性。调度器会检查：
- SM（流多处理器）上还有空闲的计算单元吗？
- 共享内存够用吗？
- 寄存器够用吗？
- 有没有数据依赖？
如果两个流里的kernel都满足条件，调度器就会让它们同时执行。否则，一个kernel会阻塞，直到资源释放。
我的经验：我曾经调试过一个程序，两个流里的kernel都用了大量共享内存。结果它们没法同时运行，因为SM上的共享内存被第一个kernel占满了。解决方案是把共享内存使用量降下来，或者让两个kernel错开执行。
4.4 流调度的优先级与抢占
你可能不知道，CUDA流是有优先级的。从CUDA 8.0开始，我们可以创建高优先级和低优先级的流。
// 创建高优先级流
int priority_high;
cudaDeviceGetStreamPriorityRange(NULL, &priority_high);
cudaStream_t high_prio_stream;
cudaStreamCreateWithPriority(&high_prio_stream,
cudaStreamNonBlocking,
priority_high);
// 创建低优先级流
int priority_low = 0;
cudaStream_t low_prio_stream;
cudaStreamCreateWithPriority(&low_prio_stream,
cudaStreamNonBlocking,
priority_low);高优先级的流会优先获得硬件资源。但注意——这不是绝对的抢占。GPU不支持像CPU那样的任务抢占。一旦一个warp开始执行，它就会一直运行到完成或者遇到内存访问延迟。
避坑指南：我曾经在实时渲染项目里用了高优先级流，以为能保证低延迟。结果发现高优先级流里的kernel太大，占用了所有SM资源，导致其他流完全无法执行。后来我把大kernel拆成多个小kernel，效果才好起来。
4.5 流调度的可视化：硬件队列状态图
为了让你更直观地理解，我画了一张流调度的工作流程图：
4.6 实际案例：流调度对性能的影响
我拿一个实际项目来说。之前做视频处理，需要同时做两件事：从CPU拷贝数据到GPU，然后在GPU上做滤波处理。
一开始我这么写：
// 单流版本 - 串行执行
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
filter_kernel<<<grid, block>>>(d_data, d_result);
cudaMemcpy(h_result, d_result, size, cudaMemcpyDeviceToHost);这个版本，拷贝和计算是串行的。总时间 = 拷贝时间 + 计算时间 + 回拷时间。
后来改成双流：
// 双流版本 - 流水线并行
cudaStream_t stream_copy, stream_compute;
cudaStreamCreate(&stream_copy);
cudaStreamCreate(&stream_compute);
// 流1：只做拷贝
cudaMemcpyAsync(d_data, h_data, size,
cudaMemcpyHostToDevice, stream_copy);
// 流2：只做计算
filter_kernel<<<grid, block, 0, stream_compute>>>(d_data, d_result);
// 流1：回拷
cudaMemcpyAsync(h_result, d_result, size,
cudaMemcpyDeviceToHost, stream_copy);你猜怎么着？性能提升了40%！因为拷贝和计算可以同时进行。硬件调度器发现拷贝队列和计算队列没有资源冲突，就让它们并行跑了。
核心结论：流调度的本质是让不同类型的任务（计算、拷贝、同步）在硬件队列上并行执行。但要注意——同类型的任务（比如两个计算kernel）能不能并行，取决于SM上是否有空闲资源。
4.7 流调度的限制与注意事项
说了这么多好处，也得说说限制。我踩过的坑不少，给你列几个：
- 资源竞争：两个流里的kernel如果都用了大量寄存器，它们就没法同时运行。因为SM上的寄存器总数是固定的。
- 调度延迟：硬件调度器切换流是有开销的。如果kernel太小（比如执行时间小于1微秒），调度开销可能比执行时间还大。
- 依赖链：如果流之间有事件同步（cudaEventSynchronize），调度器必须等待依赖满足才能继续。
- 内存带宽：多个流同时访问全局内存，可能会造成带宽争抢。我见过一个案例，两个流同时做大量内存拷贝，结果带宽被平分了，性能反而下降。
我曾经犯过的错：在一个深度学习推理项目里，我创建了16个流，每个流处理一个batch。结果发现GPU利用率只有60%。用NVIDIA Nsight分析后，发现是共享内存冲突——每个流里的kernel都申请了48KB共享内存，而SM总共只有96KB，所以最多两个流能同时跑。后来我把共享内存使用量降到16KB，利用率直接飙到95%。
嗯，总结一下。流调度机制说白了就是GPU硬件如何高效地分配资源给多个任务。它依赖Hyper-Q技术、硬件队列、以及智能的调度决策。理解了这个机制，你才能写出真正高效的并行代码。
记住一句话：流是软件，队列是硬件，调度是艺术。
公众号:蓝海资料掘金营，微信deep3321