warp 与调度

block 是给程序员看的逻辑分组,而硬件真正调度和执行的单位是 warp——32 个 lane 一组,严格锁步。

每个 block 进了 SM 之后会被切成多个 warp:

blockDim.x = 128
  → 4 个 warp(0..31, 32..63, 64..95, 96..127)

blockDim.x = 257
  → 9 个 warp(0..31, ..., 256..256 这一个 warp 只有 1 个 active lane,其它 31 个空转)

所以 blockDim 应该是 32 的倍数——否则最后一个 warp 浪费,通常 128 或 256 是最常见选择。

1. SIMT:锁步执行

warp 内 32 个 lane 在同一周期执行同一条指令。这是 NVIDIA 叫”SIMT”(Single Instruction, Multiple Thread)的来源——跟传统 SIMD(Single Instruction Multiple Data,如 AVX)在思想上一致,但调度粒度不同。

float v = a[tid] + b[tid];     // warp 内 32 个 lane 同时执行这条加法,只是各自的 tid 不同

锁步的代价是 分支分歧(warp divergence):

if (tid % 2 == 0) {
    work_A();                   // 偶数 lane 跑,奇数 lane 空转
} else {
    work_B();                   // 反过来
}

总时间 = work_A + work_B,性能减半。warp 分歧是 CUDA 优化必避的坑(详见性能优化章)。

2. SM 内部:warp scheduler

一个 SM 同时容纳几十个 warp(A100 上限是 64 个 warp / SM)。SM 内有 warp scheduler,每个 cycle 选一个 ready warp 发射指令:

SM 内部某一刻的状态:
  warp 0:正在等 global memory 数据(blocked)
  warp 1:正在等 __syncthreads(blocked)
  warp 2:有 ready 指令     ← scheduler 选它,发射
  warp 3:有 ready 指令
  ...
  warp 63

关键:warp 切换是零开销的——每个 warp 的寄存器都常驻在 SM 的寄存器文件里(不像 CPU 切换线程要存栈寄存器)。

3. 延迟隐藏:多 warp 是性能的引擎

Global memory 一次访问 ~500 cycle。如果一个 warp 一直阻塞等内存,SM 就闲下来了。

解决:让 SM 上同时有很多 warp。一个 warp 等内存时,scheduler 切到下一个 ready warp 继续算——多个 warp 互相填补对方的等待时间。

时间线:
  warp 0:[算 3 cycle][等 mem 500 cycle ............]
  warp 1:           [算 3 cycle][等 mem 500 ........]
  warp 2:                      [算 3 cycle][等 mem .]
  ...
SM 利用率:基本时刻总有一个 warp 在算

这就是为什么追求高 occupancy(SM 上活跃 warp 数 ÷ 上限)对 memory bound kernel 很重要——warp 多了才填得满 latency。

4. lane ID / warp ID

在 kernel 里查自己是哪个 warp / lane:

int tid = threadIdx.x;
int warp_id = tid / 32;          // 这个 thread 在 block 里属于第几个 warp
int lane_id = tid % 32;          // 在 warp 里是第几个 lane(0..31)

或者用更快的 %laneid PTX 特殊寄存器:

int lane_id;
asm("mov.u32 %0, %%laneid;" : "=r"(lane_id));

这两个值在写 warp-level 优化(__shfl_*、__ballot_*)时必须用。

5. warp 级原语:同步 + 通信

warp 内 32 个 lane 因为锁步,可以做一些极其便宜的协作——不需要走 shared memory:

// warp 内归约:每个 lane 加上 lane (i ^ 16) 的值,然后 ^8, ^4, ^2, ^1
for (int off = 16; off > 0; off >>= 1)
    val += __shfl_xor_sync(0xffffffff, val, off);
// 5 步之后,warp 内每个 lane 都有总和

原语	干什么
__shfl_*_sync	warp 内寄存器互换
__ballot_sync	收集 warp 内 32 个谓词成一个 uint32
__any_sync / __all_sync	warp 内”有任一/全部”为真
__syncwarp	warp 内屏障(__syncthreads 的轻量版)

这些原语不需要 shared memory + 不需要 __syncthreads,只在 warp 内 32 个 lane 之间走。是写高性能 reduce / scan / vote 的基础。

6. 一个直观的对比

同步级别	谁能用	延迟
__shfl_*_sync(warp 内)	warp 内 32 lane	~1 cycle
__syncthreads(block 内)	block 内所有 thread	~20 cycle
启动新 kernel(grid 同步)	整个 grid	~10 us(微秒)

能用 warp 级就用 warp 级——快几十倍。

7. SM 资源决定能塞多少 warp

每个 SM 资源固定。谁用得多,能同时跑的 warp 数就少(occupancy 降):

资源(A100)	总量	谁占
寄存器文件	65536 × 32-bit	每 thread 用多少寄存器
Shared memory	164 KB	每 block 用多少 __shared__
最大活跃 warp	64	硬限
最大活跃 block	32	硬限

实际 occupancy = min(寄存器限制, shared mem 限制, 上限) ÷ 最大活跃 warp 数。

8. 一句话

block 是给程序员看的,warp 是硬件实际调度的——32 个 lane 锁步执行同一条指令。SM 同时跑很多 warp 来填满 memory 延迟,这就是 GPU 隐藏延迟的核心机制。所以 blockDim 用 32 的倍数 + 避免 warp divergence + 多 warp 喂 SM 是写快 kernel 的三条底线。

下一篇:每个线程怎么知道自己处理哪个数据——索引与边界。