四、线程束与执行模型：Warp概念、Warp调度、分支发散、Warp Shuffle指令

好，咱们今天聊点硬核的。线程束（Warp）这个概念，说白了就是GPU执行指令的最小单位。你写代码时觉得是在操作一个个独立的线程，但硬件层面根本不是这么回事。我刚开始接触CUDA时也犯过这个错——以为32个线程就是32个独立的小人儿在干活，后来才发现，它们其实是绑在一根绳上的蚂蚱。

4.1 Warp是什么？为什么是32？

一个Warp就是32个线程的集合。这32个线程共享同一个指令地址，也就是说，它们在同一个时刻执行同一条指令。为什么是32？这是NVIDIA多年工程实践的结果。我记得在Fermi架构时代，这个数字就定下来了，一直沿用到现在。

举个例子，你启动了一个kernel，配置了128个线程。硬件会把这128个线程分成4个Warp：Warp0（线程0-31）、Warp1（线程32-63）、Warp2（线程64-95）、Warp3（线程96-127）。每个Warp独立地被调度到CUDA Core上执行。

4.2 Warp调度：硬件如何让Warp跑起来？

GPU里有个叫Warp Scheduler的硬件单元，它的工作就是不停地选一个Warp扔给执行单元去跑。每个SM（流多处理器）里有多个Warp Scheduler，具体数量取决于架构。比如Volta架构每个SM有4个，Turing也是4个，Ampere到了5个。

调度策略很简单：只要有可执行的Warp，就选一个出来执行。那什么时候Warp不可执行？比如在等内存访问返回、等同步指令、或者被分支卡住了。

Warp调度的核心指标是延迟隐藏。GPU用大量Warp来隐藏内存访问延迟。当一个Warp在等数据时，调度器立刻切换到另一个Warp。只要Warp数量够多，计算单元就能一直满负荷运转。

架构	每SM Warp Scheduler数	每Scheduler Warp数	最大Warp数/SM
Kepler (K80)	4	16	64
Maxwell (M40)	4	16	64
Volta (V100)	4	16	64
Turing (T4)	4	16	64
Ampere (A100)	5	16	80

4.3 分支发散：Warp的致命弱点

这是CUDA优化里最常踩的坑。因为Warp内所有线程执行同一条指令，所以当遇到if-else分支时，硬件会怎么做？

答案是：串行化。如果Warp里有些线程走if分支，有些走else分支，硬件会先执行if分支（此时走else的线程被屏蔽），再执行else分支（此时走if的线程被屏蔽）。两个分支都执行完了，Warp才继续往下走。

分支发散的影响有多大？来看个例子：

// 糟糕的写法：Warp内分支发散
__global__ void bad_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (data[idx] > 0.0f) {
            data[idx] = sqrt(data[idx]);  // 只有部分线程执行
        } else {
            data[idx] = 0.0f;             // 另一部分线程执行
        }
    }
}

// 好的写法：尽量让Warp内线程走相同路径
__global__ void good_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        // 用三元运算符避免分支
        data[idx] = data[idx] > 0.0f ? sqrt(data[idx]) : 0.0f;
    }
}

为什么三元运算符能避免分支发散？因为三元运算符在硬件层面会被编译成sel（select）指令，它不产生分支，所有线程都执行同样的指令，只是根据条件选择结果。

4.4 Warp Shuffle指令：Warp内的数据交换

Shuffle指令是我个人非常喜欢的一个特性。它允许Warp内的线程直接交换数据，不需要经过共享内存或全局内存。这在很多场景下能大幅提升性能。

Shuffle有四种变体：

• __shfl_sync()：从指定线程获取数据
• __shfl_up_sync()：从编号更小的线程获取数据
• __shfl_down_sync()：从编号更大的线程获取数据
• __shfl_xor_sync()：按XOR模式交换数据

来看个实际例子：Warp内求最大值

__device__ float warp_reduce_max(float val) {
    unsigned int mask = 0xffffffff;  // 整个Warp参与
    
    // 第一次：线程0-15从线程16-31获取数据
    val = max(val, __shfl_xor_sync(mask, val, 16));
    // 第二次：线程0-7从线程8-15获取数据
    val = max(val, __shfl_xor_sync(mask, val, 8));
    // 第三次：线程0-3从线程4-7获取数据
    val = max(val, __shfl_xor_sync(mask, val, 4));
    // 第四次：线程0-1从线程2-3获取数据
    val = max(val, __shfl_xor_sync(mask, val, 2));
    // 第五次：线程0从线程1获取数据
    val = max(val, __shfl_xor_sync(mask, val, 1));
    
    return val;  // 所有线程都得到最大值
}

Shuffle指令的延迟大约是5个时钟周期，而共享内存的延迟是20-30个周期。你想想看，这差距有多大。所以能用Shuffle的地方，千万别用共享内存。

4.5 知识体系总览

下面这张图把本章的核心逻辑串起来了：

这张图把Warp的三个核心知识点串在了一起。你想想看，Warp调度决定了你的代码能不能充分利用硬件资源，分支发散决定了你的Warp会不会空转，而Shuffle指令则提供了Warp内最高效的数据交换方式。这三者相辅相成，理解了它们，你就掌握了GPU执行模型的核心。

嗯，这一章的内容就到这里。记住：写CUDA代码时，时刻想着Warp。你的每一个线程都在一个Warp里，Warp的行为决定了你的性能上限。