一、任务调度基础：GPU任务调度模型、Stream与Concurrent Kernel、Grid/Block/Warp调度

各位同学，今天咱们聊聊GPU任务调度。说实话，我刚接触CUDA那会儿，以为把数据扔给GPU就完事了。结果呢？性能惨不忍睹。后来才明白——不理解调度模型，你连GPU的门都没摸到。

1.1 GPU任务调度模型

GPU的任务调度，说白了就是「谁先上、谁后上、谁和谁一起上」。CPU那边是操作系统管线程调度，GPU这边呢？硬件自己有一套机制。

我习惯把GPU调度模型分成三层：

• 任务级调度：Kernel函数怎么排队、怎么并行
• 线程块级调度：Block怎么分配到SM上
• 线程束级调度：Warp怎么在SM内部执行

这三层是递进关系。你想想看，一个Kernel提交上去，先得等前面的任务跑完（或者并行跑），然后它的Block被分到各个SM，最后SM里的Warp轮流执行。

1.2 Stream与Concurrent Kernel

Stream是什么？我给它一个通俗的定义：GPU上的任务队列。

默认情况下，所有操作都在同一个Stream（默认流）里执行。这意味着Kernel A没跑完，Kernel B就得等着。但如果你创建多个Stream，情况就不一样了。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 两个Kernel在不同Stream中提交
kernel_A<<<grid, block, 0, stream1>>>(...);
kernel_B<<<grid, block, 0, stream2>>>(...);

cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

这段代码里，kernel_A和kernel_B有可能同时执行——前提是硬件资源够用。我在项目中遇到过这样的情况：用4个Stream做数据流水线，吞吐量直接翻倍。但要注意，不是Stream越多越好，SM数量有限，Stream太多反而增加调度开销。

Concurrent Kernel（并发内核）是Stream的进阶玩法。它允许不同Kernel在同一设备上同时执行。嗯，这里要注意：并发执行需要满足两个条件：

• Kernel之间没有数据依赖
• 硬件资源（SM、寄存器、共享内存）足够

我建议你在Nsight Systems里看看时间轴，如果发现Kernel之间有空隙，那就是并发没做好。

1.3 Grid/Block/Warp调度

这部分是GPU调度的核心。咱们一层层拆开看。

Grid调度

Grid就是一次Kernel调用产生的所有线程块集合。GPU把Grid里的Block分配到各个SM上。分配策略是轮询（Round-Robin）吗？其实不是。硬件会尽量把Block均匀分布，同时考虑SM的负载。

Block调度

Block到了SM上，SM会为它分配资源：寄存器、共享内存、线程槽位。如果资源不够，Block就得排队等着。我记得有一次调试，发现一个SM上只跑了2个Block，利用率才30%。后来把Block尺寸调小，一个SM能塞8个Block，性能直接起飞。

Warp调度

Warp是32个线程的集合，也是SM调度的基本单位。SM里有多个Warp Scheduler，每个Scheduler负责管理一组Warp。当一个Warp在等待内存访问时，Scheduler会切换到另一个就绪的Warp——这就是延迟隐藏的秘诀。

你想想看，如果SM里只有4个Warp，那内存延迟来了，没别的Warp可切换，SM只能干等。如果SM里有32个Warp，那切换空间就大多了。

1.4 调度模型的可视化

下面这张图是我自己画的调度模型框架，帮你理清思路：

1.5 实际项目中的调度优化

最后分享一个我实际项目中的案例。当时做视频处理，需要同时做解码、缩放、编码三个Kernel。一开始全放默认流，串行执行，帧率只有30fps。

后来改成三个Stream：

1. Stream 0：解码Kernel
2. Stream 1：缩放Kernel（依赖解码结果，但通过事件同步）
3. Stream 2：编码Kernel（依赖缩放结果）

你猜怎么着？帧率从30fps飙到了85fps。原因很简单：解码和编码可以部分重叠执行，中间缩放占用的时间被隐藏了。

嗯，任务调度这块内容就讲到这里。记住一句话：调度不是玄学，是资源管理。把Stream、Block、Warp这三层理解透了，GPU性能优化就成功了一半。