3、多Stream并发调度：多Stream并行执行、硬件资源竞争、Grid Size与Block Size对并发的影响、Occupancy计算
好，咱们进入第三章。多Stream并发调度，这名字听着挺唬人，说白了就是让GPU同时干几件不同的活。我刚开始接触CUDA时，以为只要把任务扔给GPU，它就会自动并行。后来发现，嗯，没那么简单。
3.1 多Stream并行执行：别让GPU闲着
Stream，你可以把它想象成一条流水线。默认情况下，所有操作都在同一个Stream里排队执行。这就像只有一个收银台，顾客再多也得一个一个来。
多Stream就是多开几个收银台。每个Stream可以独立地执行一系列操作，包括内存拷贝和Kernel启动。这样一来，计算和传输就能重叠起来，GPU的利用率自然就上去了。
核心思想： 让GPU始终有活干，别闲着。一个Stream在等数据，另一个Stream可以算数据。
我个人习惯，在需要处理大量独立数据块时，比如视频帧处理、批量图片推理，一定会用多Stream。我曾经在一个视频分析项目中，单Stream处理时GPU利用率只有40%，改成4个Stream后，直接飙到85%以上。
3.2 硬件资源竞争：蛋糕就那么大
多Stream并行，听起来很美好，但硬件资源是有限的。GPU就那么点SM（Streaming Multiprocessor），就那么点显存带宽，还有寄存器、共享内存，这些都是共享资源。
当多个Stream同时运行时，它们会争夺这些资源。这就像几个厨师共用一间厨房，灶台、锅碗瓢盆都得抢着用。如果资源竞争太激烈，反而会降低整体性能。
| 资源类型 | 竞争影响 | 我的经验 |
|---|---|---|
| SM计算单元 | Kernel执行时间变长 | Stream太多，每个Kernel分到的SM就少 |
| 显存带宽 | 内存访问延迟增加 | 多个Stream同时读写，带宽容易打满 |
| 共享内存/寄存器 | Occupancy下降 | 每个Block占的资源多，能同时跑的Block就少 |
注意： 不是Stream越多越好。我曾经试过在一个小模型上开32个Stream，结果性能反而比4个Stream还差。资源竞争太严重了。
3.3 Grid Size与Block Size对并发的影响
Grid Size和Block Size，这两个参数直接决定了你的Kernel能产生多少线程块（Block）和线程（Thread）。它们对并发的影响，主要体现在两个方面：
- Block数量： 决定了有多少个Block可以同时被调度到SM上。Block太少，SM吃不饱；Block太多，调度开销会变大。
- Block大小： 决定了每个Block里的线程数。Block太小，SM的利用率低；Block太大，可能会受限于寄存器或共享内存，导致Occupancy下降。
你想想看，如果Grid Size只有8，而GPU有80个SM，那大部分SM都闲着。反过来，如果Block Size是1024，但每个线程用了太多寄存器，那一个SM可能只能同时跑一个Block，并发度也上不去。
我的建议： 先保证Block数量足够多（至少是SM数量的几倍），再调整Block Size。一般从128或256开始试，然后看Occupancy。
3.4 Occupancy计算：到底用了多少硬件资源？
Occupancy，直译就是「占用率」。它衡量的是：一个SM上实际活跃的Warp数量，与理论上能支持的最大Warp数量的比值。
说白了，就是看你的Kernel把SM的「坑位」占满了没有。Occupancy越高，说明SM的利用率越高，隐藏延迟的能力也越强。
计算公式其实不复杂：
Occupancy = (每个SM活跃的Warp数) / (每个SM最大支持的Warp数)但实际计算时，要考虑几个限制因素：
- 寄存器限制： 每个线程用的寄存器太多，能同时跑的Block就少。
- 共享内存限制： 每个Block用的共享内存太多，同理。
- Block大小限制： Block太小，Warp数量上不去。
- Block数量限制： 每个SM能同时跑的Block数量有上限（通常8-32个）。
举个例子，假设一个SM最多支持64个Warp，每个Block有256个线程（8个Warp）。如果每个SM最多能同时跑8个Block，那理论上最多就是8*8=64个Warp，Occupancy就是100%。
但如果每个线程用了32个寄存器，而SM总共只有65536个寄存器，那每个SM能同时跑的线程数就变成了65536/32=2048个线程，也就是8个Block。如果Block大小是512，那只能跑4个Block，Occupancy就掉到50%了。
避坑指南： 我曾经为了追求极致的单线程性能，给每个线程分配了大量寄存器。结果Occupancy掉到30%以下，整体性能反而下降了。后来我学乖了，用__launch_bounds__限制寄存器使用，Occupancy上去了，性能也回来了。
计算Occupancy，可以用NVIDIA的cuda-occupancy-calculator工具，或者在代码里用cudaOccupancyMaxActiveBlocksPerMultiprocessor这个API。我个人习惯在开发阶段，先估算一下，然后用工具验证。
嗯，这张图把整个流程串起来了。主机分发任务到多个Stream，Stream们去竞争GPU的硬件资源，而Occupancy就是衡量这个竞争效率的指标。
最后总结一下：多Stream并发调度，核心就是平衡。平衡Stream数量、Block大小、资源占用，才能让GPU跑得又快又稳。别贪多，也别太少，找到那个「甜点」位置。
实用技巧： 用nvidia-smi或者Nsight Systems监控GPU利用率。如果利用率低于70%，可以考虑增加Stream或调整Block大小。如果超过95%，那说明已经压榨得差不多了。
想要更多资料，联系可可