3、流调度与硬件：GPU硬件队列与流调度、Hyper-Q技术详解
好，我们接着聊流。上一节我们把流的概念讲清楚了，也说了怎么用。但你可能会有个疑问：我创建了这么多流，GPU到底是怎么调度它们的？ 是排队一个个来，还是能并行跑？
这就要深入到硬件层面了。说白了，流调度这件事，硬件和软件是两码事。你软件上分了32个流，硬件不一定能同时处理32个。这里面的门道，我今天给你讲透。
3.1 硬件队列：GPU的“任务池”
每个GPU内部，都有一组硬件工作队列（Hardware Work Queue）。你可以把它想象成一个“任务池”。
CPU提交任务时，其实是把任务扔到这个池子里。GPU的调度器再从池子里捞任务出来执行。
在Fermi架构之前，情况很简单：只有一个硬件队列。所有流里的任务，最终都挤到这一个队列里。这就带来了一个问题——
伪依赖（False Dependency）：明明两个任务没有数据依赖，但因为挤在同一个队列里，必须等前一个执行完，后一个才能开始。
我记得早期做项目时，用CUDA 2.x的卡，明明分了4个流，结果性能提升微乎其微。查了半天，发现是硬件队列不够用。所有流都在排队，跟单流没区别。
嗯，这里要注意：软件流再多，硬件队列不够，也是白搭。
3.2 从Kepler到Hyper-Q：一次革命
到了Kepler架构（K20、K40那代），NVIDIA终于开窍了。他们引入了Hyper-Q技术。
Hyper-Q是什么？说白了，就是把硬件队列从1个，变成了32个。
你想想看，32个硬件队列，对应32个软件流。每个流都有自己的专属通道。任务提交时，直接扔到对应的硬件队列里。调度器可以同时从32个队列里捞任务，真正的并行执行就来了。
我当年第一次用K40跑多流测试时，看到NVIDIA Visual Profiler里32个队列同时活跃，那种感觉——嗯，就像堵车时突然多了32条车道。
个人经验：Hyper-Q对CPU端的多线程提交特别友好。如果你有多个CPU线程各自提交任务到不同流，Hyper-Q能保证每个线程的任务不会被其他线程阻塞。我曾经在一个实时渲染项目里，用4个CPU线程分别提交计算、拷贝、渲染任务，配合Hyper-Q，延迟降低了40%。
3.3 流调度的硬件机制
有了Hyper-Q，调度器怎么工作？
GPU内部有一个调度单元（Warp Scheduler），它负责从硬件队列里取任务。取任务的策略是：
- 轮询（Round-Robin）：每个队列轮流看一眼，有任务就取一个。
- 优先级：某些队列可以设置高优先级，调度器会优先处理。
- 依赖检查：取任务前，检查该任务依赖的事件（Event）是否已经完成。
这里有个关键点：同一个流内的任务，必须按顺序执行。这是CUDA的语义保证。但不同流的任务，调度器可以乱序执行。
举个例子：
// 流A
cudaStream_t streamA, streamB;
cudaStreamCreate(&streamA);
cudaStreamCreate(&streamB);
// 流A提交任务1
kernelA1<<<grid, block, 0, streamA>>>(...);
// 流B提交任务2
kernelB1<<<grid, block, 0, streamB>>>(...);
// 流A提交任务3
kernelA2<<<grid, block, 0, streamA>>>(...);硬件调度时，可能先执行kernelB1，再执行kernelA1，再执行kernelA2。但kernelA2一定在kernelA1之后，因为它们在同一个流里。
避坑指南：我曾经遇到过一个bug，在流A里先提交了cudaMemcpyAsync，再提交kernel。结果kernel先执行了，因为cudaMemcpyAsync还在队列里没被调度到。后来发现是没有正确使用事件同步。记住：同一个流内虽然顺序保证，但如果你用多个流，依赖关系必须显式声明。
3.4 Hyper-Q的局限性
Hyper-Q虽然好，但不是万能的。有几个限制你得知道：
| 限制项 | 说明 |
|---|---|
| 队列数量上限 | Kepler及之后架构最多32个硬件队列。你创建64个流，也只有32个能真正并行。 |
| 资源竞争 | 32个队列共享SM资源。如果每个队列都提交大量任务，SM会过载，反而降低性能。 |
| 调度开销 | 队列越多，调度器的轮询开销越大。极端情况下，调度本身会成为瓶颈。 |
我个人习惯是：流数量不要超过SM数量的2-3倍。比如一个GPU有16个SM，我一般用32-48个流。再多，收益就递减了。
3.5 实际项目中的流调度策略
讲了这么多理论，来点实际的。我在做深度学习训练框架时，总结了一套流调度策略：
- 计算密集型任务：用2-4个流就够了。太多流反而增加调度开销。
- I/O密集型任务（频繁的memcpy）：用8-16个流。让拷贝和计算重叠。
- 混合负载：用16-32个流。把计算、拷贝、控制任务分开到不同流。
你可能会问：怎么知道当前GPU有多少硬件队列？
很简单，查一下计算能力（Compute Capability）。3.0及以上都支持Hyper-Q。具体数量：
- Kepler (3.0-3.7): 32个队列
- Maxwell (5.0-5.3): 32个队列
- Pascal (6.0-6.2): 32个队列
- Volta/Turing/Ampere (7.0-8.9): 32个队列
嗯，从Kepler到现在，32个队列这个数字一直没变。我猜NVIDIA觉得32个够用了。确实，大多数场景下，32个流已经能榨干GPU了。
核心总结：Hyper-Q让多流真正并行成为可能。但别盲目创建大量流。理解硬件队列的机制，合理分配任务，才能发挥最大性能。
下一节，我们会讲流与事件的高级用法，包括如何用事件做细粒度同步、如何避免死锁。到时候我会分享一个我踩过的坑——因为事件使用不当，导致整个GPU卡死。嗯，那是个难忘的教训。