<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 2 · 执行模型与内存04 / 15约 24 分钟

Warp 与 SM：GPU 真正的执行单位

SIMT 模型、warp divergence、延迟隐藏，以及「32」这个数字为什么无处不在

学完这一课你会
  • 理解 warp 是硬件调度的最小单位，而不是线程
  • 看懂 warp divergence 的代价，并能重写代码规避它
  • 明白 GPU 靠什么隐藏访存延迟——不是缓存，是超额线程
  • 区分 __syncthreads() 和 __syncwarp() 的适用场景

编程模型里你写的是「线程」，但硬件根本不以单个线程为单位调度。32 个连续的线程被绑成一个 warp，共享同一个程序计数器，锁步执行同一条指令。这就是 NVIDIA 说的 SIMT（Single Instruction, Multiple Thread）。理解 warp，是从「能写对 CUDA」到「能写快 CUDA」的分水岭。

┌──────────────────────── SM ────────────────────────┐
│  Warp Scheduler ×4   每个周期各挑一个 ready 的 warp  │
│         │                                          │
│  ┌──────┴──────┬───────────┬──────────┐            │
│  │ CUDA Core   │  LD/ST    │  SFU     │  Tensor    │
│  │  ×128       │  单元      │          │  Core ×4   │
│  └─────────────┴───────────┴──────────┴────────────┘│
│                                                     │
│  Register File   256 KB（全 SM 共享，按线程切分）      │
│  Shared Memory / L1   最多 164 KB（可配置切分）        │
└─────────────────────────────────────────────────────┘
              │
        ┌─────┴─────┐
        │  L2 Cache │  全芯片共享，几十 MB
        └─────┬─────┘
              │
        ┌─────┴─────┐
        │   HBM     │  显存，几十 GB，带宽 1~3 TB/s
        └───────────┘
一个 SM 的简化结构（Ampere 类）

一个 block 是怎么被执行的

  1. 1.调度器把整个 block 原子地分配给某一个 SM——block 不会被拆到两个 SM 上。
  2. 2.SM 按线性 tid 把 block 切成若干 warp。blockDim = 256 就是 8 个 warp。
  3. 3.如果 blockDim 不是 32 的倍数，最后一个 warp 会被填满但部分线程处于 inactive 状态，这些槽位的算力直接浪费。
  4. 4.SM 上的 warp 调度器每个周期从所有驻留 warp 里挑出「就绪」的发射指令。一个 warp 在等访存时，调度器立刻切到别的 warp——零开销切换，因为每个 warp 的寄存器一直物理保留着。

Warp Divergence：if 的真实代价

一个 warp 只有一个 PC。当 warp 里的线程在 if 上走了不同分支，硬件没法同时执行两条路径，只能串行执行两遍：先跑 then 分支（走 else 的线程被掩码屏蔽、空转），再跑 else 分支（反过来屏蔽）。两条分支的时间相加，这就是 warp divergence。

最坏的分支写法
1__global__ void bad(float* x) {2    int i = blockIdx.x * blockDim.x + threadIdx.x;3 4    if (i % 2 == 0) {         // warp 内 32 个线程一半一半，100% 分叉5        x[i] = expensiveA(x[i]);6    } else {7        x[i] = expensiveB(x[i]);8    }9    // 实际耗时 ≈ expensiveA + expensiveB，而不是二选一10}
按 warp 粒度对齐分支：不分叉
1__global__ void good(float* x) {2    int i = blockIdx.x * blockDim.x + threadIdx.x;3 4    if ((i / warpSize) % 2 == 0) {   // 同一个 warp 内所有线程走向一致5        x[i] = expensiveA(x[i]);6    } else {7        x[i] = expensiveB(x[i]);8    }9    // 每个 warp 只执行一条路径，代价回到二选一10}

同步：__syncthreads 与 __syncwarp

block 级屏障
1__global__ void stencil(const float* in, float* out, int n) {2    __shared__ float tile[BLOCK + 2 * RADIUS];3 4    int gid = blockIdx.x * blockDim.x + threadIdx.x;5    int lid = threadIdx.x + RADIUS;6 7    tile[lid] = in[gid];                       // 每个线程写自己那格8    if (threadIdx.x < RADIUS) {                // 前几个线程顺便搬 halo9        tile[lid - RADIUS]     = in[gid - RADIUS];10        tile[lid + blockDim.x] = in[gid + blockDim.x];11    }12 13    __syncthreads();     // 屏障：确保 tile 全部写完，才允许任何人开始读14 15    float sum = 0.0f;16    for (int d = -RADIUS; d <= RADIUS; ++d) sum += tile[lid + d];17    out[gid] = sum;18}

从 Volta（sm_70）开始，NVIDIA 引入了独立线程调度（Independent Thread Scheduling）：warp 内的线程各自拥有 PC，分叉后不再保证会自动重新汇聚。这带来一个后果——过去那种「同一个 warp 内天然同步，所以不用加同步」的老代码（尤其是 warp 级归约）在 Volta 之后是错的。现在所有 warp 内的数据交换都必须显式带上掩码同步。

Volta 之后的正确写法
1// ❌ Kepler 时代的老写法，Volta 之后不再保证正确2volatile float* v = sdata;3if (tid < 32) { v[tid] += v[tid + 32]; v[tid] += v[tid + 16]; /* ... */ }4 5// ✅ 显式同步版本6if (tid < 32) {7    float val = sdata[tid] + sdata[tid + 32];8    for (int offset = 16; offset > 0; offset >>= 1) {9        val += __shfl_down_sync(0xffffffff, val, offset);10    }11    if (tid == 0) out[blockIdx.x] = val;12}

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。