4、线程调度机制：Warp/Wavefront概念、两级线程调度器、线程块分配策略

好，咱们今天聊一个GPU架构里非常核心的话题——线程调度。说白了，GPU能同时处理成千上万个线程，靠的就是一套高效的调度机制。我当年刚接触GPU设计时，最困惑的就是：这么多线程，到底怎么管？谁先执行？谁后执行？

嗯，今天咱们就把这个问题彻底讲透。

4.1 Warp与Wavefront：调度的基本单位

先说说Warp。NVIDIA叫它Warp，AMD叫它Wavefront。名字不同，本质一样——都是一组同时执行的线程。

我个人习惯把Warp理解为“GPU里的最小调度单元”。一个Warp通常包含32个线程（NVIDIA）或64个线程（AMD）。这32个线程共享同一个程序计数器，执行同一条指令，只是处理的数据不同。

你可能会问：为什么是32？不是16，也不是64？

这其实跟硬件设计有关。32个线程同时执行，既能充分利用ALU资源，又不会让调度器太复杂。我在项目中遇到过，如果把Warp大小改成64，虽然理论吞吐量翻倍，但实际因为分支发散和资源冲突，性能反而下降。所以32是个很平衡的选择。

举个例子：

// 假设一个Warp有32个线程
// 每个线程处理一个数组元素
__global__ void vec_add(float* a, float* b, float* c) {
    int idx = threadIdx.x;  // 0~31
    c[idx] = a[idx] + b[idx];
}

这段代码里，一个Warp的32个线程同时执行加法。硬件上，它们共享指令，但各自有独立的寄存器。

4.2 两级线程调度器：为什么需要两级？

好，现在我们知道Warp是基本单位了。那调度器怎么工作？

GPU的调度器是两级结构。我刚开始看这个设计时觉得多此一举，后来才明白——这是为了平衡延迟隐藏和资源利用率。

第一级：线程块调度（Block-level Scheduler）

这一级负责把线程块（Thread Block）分配到SM（流多处理器）上。说白了，就是决定哪个线程块去哪个SM干活。

分配策略很简单：哪个SM有空闲资源，就分配过去。资源包括寄存器、共享内存、线程槽位等。

第二级：Warp调度（Warp-level Scheduler）

这一级才是真正的“微观调度”。每个SM内部有一个Warp调度器，它负责从当前驻留的所有Warp中，选出一个来执行。

选谁？常见策略有：

• 轮转调度（Round-Robin）：每个Warp轮流执行，公平但效率一般
• 优先级调度：给某些Warp更高优先级，适合实时任务
• 最老优先（Oldest-first）：等待时间最长的Warp先执行，延迟隐藏效果好

我个人最常用的是最老优先。为什么？因为GPU最大的敌人是内存延迟。一个Warp在等数据时，调度器可以切换到另一个Warp。最老优先能保证没有Warp被饿死。

看个调度流程：

// 伪代码：两级调度器工作流程
while (有线程块待分配) {
    // 第一级：分配线程块到SM
    SM = find_idle_sm();
    if (SM != NULL) {
        block = get_next_block();
        SM.assign(block);
    }
    
    // 第二级：每个SM内部调度Warp
    for each SM {
        warp = select_warp_oldest_first();
        warp.execute_next_instruction();
    }
}

4.3 线程块分配策略：怎么分才高效？

线程块分配看似简单，其实门道很多。我踩过不少坑，分享几个关键点。

策略一：资源感知分配

分配线程块时，必须知道SM还剩多少资源。比如：

资源类型	SM总量	每个线程块需求	最大线程块数
寄存器	65536	16384	4
共享内存	48KB	16KB	3
线程槽位	2048	512	4

你看，共享内存成了瓶颈，最多只能放3个线程块。这时候强行分配第4个，只会导致等待。

策略二：负载均衡

多个SM之间要尽量均衡。如果一个SM忙死，另一个SM闲死，那整体性能就上不去。

常见的做法是轮询分配：把线程块依次分配给SM0、SM1、SM2... 这样每个SM分到的线程块数量差不多。

但有个细节：如果线程块的计算量不均匀，轮询也没用。比如一个线程块要算1000次，另一个只算10次。这时候需要动态负载均衡——让空闲的SM去“偷”任务。

策略三：Warp级并行度

每个SM能同时驻留的Warp数量，决定了它的并行度。并行度越高，延迟隐藏能力越强。

我建议：尽量让每个SM的Warp数量达到上限。比如SM最多64个Warp，那就分配足够的线程块，让Warp数接近64。

为什么？因为内存访问延迟通常几百个周期。如果只有8个Warp，那8个Warp都等内存时，ALU就闲着。如果有64个Warp，总有几个Warp的数据准备好了，ALU就能一直干活。

4.4 实际设计中的避坑指南

最后，分享几个我亲身经历的教训：

• 别让Warp发散太严重：如果一个Warp里32个线程走不同的分支，那性能会退化到串行。我遇到过最夸张的情况，一个Warp里32个线程走了32条不同路径，性能直接掉到1/32。
• 注意线程块大小：线程块太小（比如32个线程），SM利用率低。线程块太大（比如1024个线程），资源竞争激烈。我一般推荐256~512个线程。
• 调度器本身也有开销：两级调度器虽然灵活，但每次切换Warp都要保存/恢复状态。如果Warp切换太频繁，调度开销会吃掉性能。

嗯，关于线程调度，今天就聊这么多。记住一句话：调度的本质，是用计算掩盖延迟。理解了这句话，你就抓住了GPU性能优化的核心。