任务调度基础：GPU任务调度模型、Stream与Concurrent Kernel、硬件调度单元概览
大家好，欢迎来到任务调度深度剖析的第二讲。今天咱们聊聊GPU到底是怎么把任务安排得明明白白的。说实话，我早年刚接触CUDA时，以为把数据扔给GPU就完事了，结果性能惨不忍睹。后来才明白——调度，才是GPU性能的灵魂。
一、GPU任务调度模型：从CPU到GPU的“搬运”逻辑
GPU不是独立存在的，它得听CPU的指挥。CPU把任务“发射”到GPU，这个过程叫任务提交。你想想看，CPU和GPU是两套完全不同的硬件，它们之间怎么协作？
说白了，CPU通过驱动把kernel、内存操作、同步指令打包成一个命令缓冲区，然后通过PCIe总线推给GPU。GPU这边有个前端控制器（也叫Host Interface），负责接收这些命令，再分发给后面的执行单元。
我个人习惯把GPU任务调度分成三个层次：
- 主机端调度：CPU侧的任务提交与队列管理
- 设备端调度：GPU内部的命令分发与执行
- 硬件级调度：SM内部的warp调度
嗯，这里要注意，这三个层次是层层递进的。很多初学者只关注第三个，忽略了前两个，结果就是——CPU在等GPU，GPU在等数据，两边都在空转。
二、Stream与Concurrent Kernel：并行执行的“车道”
Stream，你可以把它想象成GPU里的任务流水线。每个Stream里的操作是顺序执行的，但不同Stream之间可以并行。我遇到过不少项目，开发者把所有操作都塞进默认Stream（Stream 0），结果GPU利用率只有30%。
为什么会这样？因为默认Stream有个特殊限制——它会阻塞其他Stream。说白了，默认Stream像个“霸道总裁”，它执行时别的Stream都得等着。
关键点：要真正实现Concurrent Kernel，必须使用非默认Stream，并且确保kernel之间没有数据依赖。
来看一个实际例子：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 两个kernel可以并发执行
kernel_A<<<grid, block, 0, stream1>>>(d_data1);
kernel_B<<<grid, block, 0, stream2>>>(d_data2);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);这里有个坑——我曾经以为只要用了不同Stream就能自动并行。结果发现，如果两个kernel访问同一块显存的不同区域，硬件可能会因为内存带宽争抢而串行化。所以，Concurrent Kernel不仅要看Stream，还得看内存访问模式。
我的建议：用Nsight Systems的Timeline视图观察Stream活动。如果两个kernel的时间线没有重叠，说明它们没有真正并发。这时候检查一下：是不是用了默认Stream？是不是有隐式同步？
三、硬件调度单元：GPC/TPC/SM概览
好了，现在咱们深入到芯片内部。NVIDIA GPU的硬件调度单元，说白了就是一套分级管理体系。我以Ampere架构为例（GA102核心），给你拆解一下：
| 层级 | 全称 | 作用 | 数量（GA102） |
|---|---|---|---|
| GPC | Graphics Processing Cluster | 负责图形/计算任务的粗粒度分配 | 7个 |
| TPC | Texture Processing Cluster | 管理纹理单元和SM的配对 | 每个GPC含2个TPC |
| SM | Streaming Multiprocessor | 真正的计算执行单元 | 每个TPC含2个SM |
你想想看，一个kernel启动时，硬件是怎么找到空闲SM的？GPC里的Work Distribution Unit（工作分发单元）会扫描所有SM的负载情况，然后把线程块（Thread Block）分配给负载最轻的SM。这个过程是硬件自动完成的，我们开发者没法直接干预。
但我可以告诉你一个小技巧——通过控制线程块大小和共享内存使用量，可以间接影响SM的占用率。我做过一个实验：把线程块从256改成128，虽然线程总数不变，但SM能同时容纳更多线程块，延迟隐藏效果更好。
注意：SM内部的warp调度器才是真正的“微观调度者”。每个SM有4个warp调度器（Ampere架构），每个调度器每周期可以发射一条指令。如果warp因为访存或同步而停顿，调度器会立刻切换到另一个就绪的warp——这就是GPU延迟隐藏的核心机制。
四、知识体系总览
下面这张图是我自己整理的，把今天讲的内容串起来。你看一眼就能明白整个调度链条：
这张图里，从CPU发射任务开始，经过命令缓冲区、GPC、TPC，最后落到SM上执行。每个环节都有对应的硬件单元负责。我当年调优时，就是盯着这张图，一个一个环节排查瓶颈。
五、避坑指南与实战心得
最后分享几个我踩过的坑：
- Stream数量不是越多越好：我曾经一口气创建了32个Stream，结果性能反而下降了。因为硬件能同时处理的Stream数量有限（通常16个左右），太多Stream会导致调度开销增加。
- Concurrent Kernel需要硬件支持：不是所有GPU都支持。检查一下你的设备属性——
cudaDeviceProp.concurrentKernels。我遇到过在旧款Tesla卡上怎么调都不并行的情况，一查发现硬件不支持。 - 内存操作也会占用Stream：
cudaMemcpyAsync同样会排队到Stream里。如果你在Stream里先做memcpy再做kernel，它们不会并行——除非你用两个不同的Stream分别处理。
一句话总结：任务调度的本质是让GPU始终有活干。Stream是软件层面的并行手段，GPC/TPC/SM是硬件层面的执行单元。理解这两者的配合，你就能写出真正高效的CUDA程序。
好了，这一讲就到这里。记住，调度不是玄学，是工程。多动手用Nsight抓一抓时间线，你会有更直观的感受。
想要更多资料，联系可可