GPU 执行模型 2：GPU 后端

thread 的定义：PC 加寄存器

不管在哪个语境下，thread 的本质都是：一段有自己的程序计数器（PC）和寄存器状态、能独立往前执行指令的控制流。

关键词是"独立的 PC + 独立的寄存器"。只要满足这个，就可以叫 thread。至于这个 thread 是被操作系统调度的、被硬件直接执行的、还是被 GPU 一组一组打包执行的，那是实现层面的差别。

CUDA 里一个 thread 是最细粒度的执行单位——你写的 kernel 函数，每个 thread 都会执行一遍，只不过通过 threadIdx 知道自己是第几个，从而处理不同的数据。

但是！CUDA 的 thread 不符合传统 thread 的定义，虽然每个 thread 有独立的寄存器，但是同一个 warp 里面的 thread，共享一个 PC。

真正符合传统 thread 定义的，其实是 warp。不同的 warp 有独立的 PC 和独立的寄存器。

那为什么 NVIDIA 还要叫每个 lane 是"thread"？因为从编程模型看，每个 lane 有自己的寄存器、自己的局部变量、可以走不同的控制流分支（虽然分支发散会损失性能），程序员写代码时可以把它当成一个独立的执行序列来思考。这是一种刻意的抽象：把 SIMD 包装成"很多个 thread"，让程序员更容易写并行代码。NVIDIA 给这套东西起了个名字叫 SIMT（Single Instruction, Multiple Threads），就是强调这种"硬件是 SIMD，但编程模型是 thread"的混合。

kernel 到硬件的映射

阶段 0：launch 配置进入 GTE

当 host 调用 kernel<<<gridDim, blockDim, smem, stream>>>(...) 时，driver 把这次 launch 打包成一个命令写进 host 上对应 stream 的命令缓冲（push buffer），最终通过 GPU 的前端（host interface / front end）送到 GPU 的全局调度硬件 GigaThread Engine。

GigaThread Engine 拿到的信息包括：

这些资源数字非常关键，因为下一步是资源能不能塞下的判断。

阶段 1：GTE 把 block 分发到 SM

GigaThread Engine 是一个 block 级别的分发器。它维护一个全局的"待发射 block 队列"，并不断尝试把 block 派到有空间的 SM 上。

注意它派的单位是 block，而不是 thread 或 warp——这是 CUDA 编程模型的核心约束之一：一个 block 一旦开始执行，必须整体待在同一个 SM 上直到结束，因为 block 内的 shared memory 和 __syncthreads() 都依赖这一点。

1a. GTE 如何判断一个 SM 能不能接收一个 block

GTE 如何判断一个 SM 能不能接收一个 block，主要看以下四个资源是否同时够用：

只要其中任何一个资源不够，这个 SM 就不能再接 block，GigaThread 会试别的 SM；如果所有 SM 都不够，新 block 就在队列里等，直到某个 SM 上有 block 退出释放资源。这就是为什么 occupancy（实际驻留 warp / 理论最大 warp）由"最紧的那个资源"决定——常见的瓶颈是寄存器。

1b. 多个 block 到多个 SM 的派发策略

1c. Wave 和 Tail Effect

1d. Concurrent Kernel Execution 和 Grid Slot

阶段 2：block 进 SM 后的资源绑定

一个 block 被派到某个 SM 上，SM 内部要做几件事，这些都是硬件在 block 启动时一次性完成的静态绑定：

2a. 把 block 内的 thread 切成 warp

按 threadIdx 的线性顺序（先 x、再 y、再 z）每 32 个 thread 打包成一个 warp。比如 blockDim = (64, 2, 1) = 128 thread → 4 个 warp，warp 0 = thread 0–31，warp 1 = thread 32–63，依此类推。这个切分是纯几何的、确定性的，硬件直接按公式算。

2b. 初始化 block slot 状态

每个 block 拿到 1 个 block slot，里面包含：

2c. 把 warp 分到 4 个 partition

GA10x 的 SM 有 4 个 partition（processing block），每个 partition 有自己独立的 warp scheduler、寄存器堆、L0 i-cache、INT/FP/Tensor 单元。

block 内的 warp 会被分散到 4 个 partition 上，常见做法是按 warp ID mod 4。这样 block 里的 warp 0、4、8 进 partition 0，warp 1、5、9 进 partition 1，等等。

为什么要分散？因为每个 partition 是一个独立的执行引擎，分散开才能并行起来。如果一个 block 只有 1 个 warp，那只有 1 个 partition 在干活，其他 3 个对这个 block 而言是闲置的——这就是为什么 block 太小（< 128 thread）通常浪费 SM 资源。

2d. 在 partition 的寄存器堆里给 warp 划寄存器

每个 partition 的 16384 个寄存器是物理上连续的一大块 SRAM。block 启动时，硬件给落到这个 partition 的每个 warp 划一段连续区域，大小 = 每 thread 寄存器数 × 32。

比如 kernel 用 40 个寄存器/thread，那每个 warp 占 40 × 32 = 1280 个寄存器，这个 partition 最多能放 16384 / 1280 ≈ 12 个 warp（正好是 GA10x 每 partition 的 warp slot 上限，所以 40 寄存器/thread 是 GA10x 上达到 100% occupancy 的临界点）。

划好之后，这个 warp 在它生命周期里访问寄存器 R5 时，硬件就知道是去寄存器堆里它那块的偏移 5 处取——这个映射通过 per-warp base register 之类的机制实现，warp 内的寄存器编号是相对的。映射信息记录在 warp slot 内。

2e. 分配 shared memory

block 整体在 SM 的 shared memory 里拿到一块连续区域，所有同 block 的 warp 共享访问。shared memory 是 per-block 的，不是 per-warp 也不是 per-partition——尽管物理上 shared memory 是分 bank 的（GA10x 是 32 bank），但地址空间对整个 block 可见。

2f. 初始化 warp slot 状态

每个 warp 拿到一个 partition 内的 warp slot，里面包含：

到这一步，block 就"就绪"了，它的所有 warp 都已经进入各自 partition 的待调度池。

阶段 3：warp scheduler 的每周期发射

现在进入稳态运行。每个 partition 的 warp scheduler 每个 cycle 做一件事：

GA10x 每个 partition 每 cycle 只能发射 1 条 warp 指令。这条指令在执行单元里以 SIMD 方式进入 lane。

关键点：这里没有任何"上下文切换"的开销。warp A 的指令发射后下个 cycle 立刻发 warp B 的指令，因为 A 和 B 的 PC、寄存器物理上同时存在于 partition 的硬件里，scheduler 只是选 mux 走哪一路。这就是 GPU 靠超大规模 SMT 隐藏延迟的物理基础——只要驻留 warp 足够多，每个 cycle 总能找到一个 ready 的，执行单元就不会闲。

另外需要注意，编译器在编译时会重排指令顺序，把独立的指令尽量提前、把有依赖的指令尽量推后。这样子，scheduler 可以尽可能地一直发射指令。

阶段 4：block 退出与资源回收

当一个 block 内所有 warp 都执行到 kernel 结尾（ret）后，等待所有 in-flight 访存落地，硬件标记这个 block 完成，回收它占用的：

回收后 SM 通知 GigaThread Engine "我这边腾出空间了"，GigaThread 就可以再派新 block 过来。整个过程就是一个生产-消费流水：grid 里上万个 block 持续往 SM 池子里灌，SM 跑完一个吐一个。

注意，warp 执行到 ret 后不释放任何资源，只是变成 not eligible 安静地占着 slot；必须等同一个 block 的所有 warp 都 ret 了，硬件才会以 block 为单位一次性回收 warp slot、寄存器、shared memory、block slot，然后通知 GigaThread Engine 派新 block 进来。

Kernel 执行时间

总结

occupancy

如何选择 blockDim 和 gridDim

Active Mask

硬件机制

warp slot 中 active mask 为 0 的 lane（thread）不会执行任何操作，不是"执行了再丢掉结果"。

硬件实现是门控（gating）：mask bit 为 0 的 lane，其对应的执行逻辑被关掉。具体到不同类型的指令：

所以 active mask 的语义是：mask bit = 0 的 lane 在这条指令上完全不存在，就好像这个 lane 对应的 thread 从来没有过一样。不发访存、不做计算、不写结果。

用途

active mask 的用途：

a. 分支发散

这是最核心的用途。遇到 if/else 时：

if (threadIdx.x < 16) { // 指令地址 100
          // 地址 101-200
    A();  // mask = 0x0000FFFF，前16个lane active
} else {
          // 地址 201-205
    B();  // mask = 0xFFFF0000，后16个lane active
}
C();      // 地址 206
          // 汇合后 mask = 0xFFFFFFFF

两个分支各自执行时，不该参与的 lane 被 mask 掉。

整个过程的关键是：硬件用栈自动管理"先跑谁、后跑谁、在哪汇合"，编译器在编译时把汇合点（RPC）算好嵌入指令里，硬件运行时机械地压栈弹栈。程序员完全不用操心。

嵌套 if 也能处理——每嵌套一层多压一层栈，发散越深栈越高。

b. 尾部 warp 不满 32 个 thread

如果 blockDim 不是 32 的倍数，最后一个 warp 的多余 lane 从 block 启动的那一刻起 mask 就是 0，整个 kernel 执行期间这些 lane 永远不会被激活。

比如 blockDim = 100，切成 4 个 warp：warp 0–2 各 32 个 thread（mask = 0xFFFFFFFF），warp 3 只有 4 个 thread（mask = 0x0000000F）。warp 3 的后 28 个 lane 全程白占着 ALU 宽度。

这就是为什么 blockDim 应该是 32 的倍数——否则你白白浪费一个 warp 的大部分执行资源，而且那些 lane 对应的寄存器也照样被分配出去了，全浪费。

c. warp 级条件执行指令

__ballot_sync、__shfl_sync、__reduce_sync 这些 warp 级原语都接受一个 mask 参数，指定哪些 lane 参与操作。硬件用这个 mask 来决定哪些 lane 贡献数据、哪些不参与。

d. 提前退出（early return / exit）

如果 kernel 里有：

if (someCondition) return;
// 后续代码

满足条件的 thread 提前退出，硬件把它们的 mask bit 清零。之后这些 lane 就和尾部 warp 的多余 lane 一样，永远不再被激活。

如果一个 warp 的所有 32 个 lane 都被清零了（全部 thread 都 return 了），这个 warp 就变成 not eligible，相当于提前 ret。

ITS

上面提到了 active mask 的一个用途，即分支发散和汇聚。但是那只是 volta 架构之前的情况，从 volta 开始的架构采用了 independent thread scheduling。

Warp 死锁

我们来看一个最经典的同 Warp 生产者-消费者（Spinlock 自旋锁）模型。

假设 Warp 里的 Thread 0 是消费者（等数据），Thread 1 是生产者（发数据）。它们通过一块 Shared Memory 里的标志位 flag 来通信。

__shared__ volatile int flag = 0;

if (threadIdx.x == 0) {
    // 消费者：死等 flag 变成 1
    while (flag == 0) {
        // 空转自旋
    }
    printf("Thread 0 收到数据！\n");
}
else if (threadIdx.x == 1) {
    // 生产者：负责把 flag 变成 1
    flag = 1;
    printf("Thread 1 发送数据！\n");
}
// 其他 30 个 thread 这里忽略

老架构（Volta 之前）会死锁。当执行到 if 的时候，发散产生了。硬件会把 else 分支 (Thread 1) 压入栈中存档，然后用 if 分支 (Thread 0) 的状态更新 Warp slot，开始执行。

此时，灾难发生了：

这就是老架构臭名昭著的同 Warp 死锁（Intra-Warp Deadlock）。它本质上是硬件强制串行化导致的饥饿现象。

ITS 如何解决

// 初始状态：32 个线程 PC 都在 100
if (threadIdx.x < 16) { // 分支指令地址 100
          // 地址 101-200
    A();  // 前 16 个 lane 需要执行
} else {
          // 地址 201-205
    B();  // 后 16 个 lane 需要执行
}
C();      // 汇合点地址 206

老架构和新架构比较

// PC 100:
if (threadIdx.x < 16) {
    A1(); // PC 101
    A2(); // PC 102
} else {
    B1(); // PC 201
    B2(); // PC 202
}
C();      // PC 300 (汇合)
执行阶段   老架构 (SIMT Stack)      新架构 (Volta+ ITS)
Cycle 1   发射 A1 (Mask: 前16)。
          后16线程在栈里吃灰。        发射 A1 (Mask: 前16)。
Cycle 2   发射 A2 (Mask: 前16)。    调度器可以根据优先级，选择发射 B1 (Mask: 后16)。
Cycle 3   if 跑完，弹栈。
          开始发射 B1 (Mask: 后16)。  调度器继续发射 A2 (Mask: 前16)。
Cycle 4   发射 B2 (Mask: 后16)。     调度器发射 B2 (Mask: 后16)。
结果       严格串行 (A1->A2->B1->B2)   指令级交替执行 (A1 -> B1 -> A2 -> B2)

ITS 并不强制必须交替执行，调度器会根据寄存器依赖、访存延迟等实际情况，灵活决定当前 cycle 挑哪个 PC 来动态打包。

ITS 的好处和注意事项

为每个 thread 增加独立的 PC 和 Call Stack 消耗了大量的 SRAM 晶体管，NVIDIA 这么做主要有两个好处：

但是需要注意，ITS 会导致一种极其危险的情况：跑得快的线程，根本不停下来等跑得慢的线程。比如：Thread 0-15 走 if 并且瞬间跑完了；Thread 16-31 走 else 并且因为去读全局内存卡住了。在 ITS 下，调度器可能让 Thread 0-15 的 PC 顺着代码直接往下冲，进入了后面的计算阶段！

如果不加 __syncwarp()，跑得快的 Thread 0 就直接往下冲，去读 Thread 16 应该写但还没来得及写的数据，直接读出一堆垃圾值。

__syncthreads

例子

用一个具体例子：block 有 8 个 warp（warp 0–7），分散在 4 个 partition 上（每个 partition 2 个）。

在 block 启动时，SM 级的 block slot 里的所有 Barrier 单元会初始化一个 Participation Mask（参与位图）。因为这个 block 有 8 个 warp，所以位图初始化为 11111111（表示期待这 8 个 warp 都来报到），同时维护一个初始为全 0 的 Arrival Mask（到达位图）。

a. 各 warp 陆续到达

假设 warp 2 最先执行到 __syncthreads() 那条指令（这在底层对应的是默认的 BAR.SYNC 0，即等待 0 号 Barrier）：

注意这里的分工：warp slot 在 partition 里负责“挂起自己并记住在等谁”，barrier 单元组在 SM 级的 block slot 里负责“全局统筹”。这是因为同一个 block 的 warp 分散在 4 个 partition 上，barrier 必须在 SM 级统一记账。

b. 更多 warp 到达

在这个过程中，还没到达 barrier 的 warp 继续正常执行——它们还是 eligible 的，scheduler 照常 issue 它们的指令。已经到达的 warp 安静地占着 warp slot 和寄存器，但不消耗 issue 带宽。

c. 最后一个 warp 到达

warp 6 是最后一个到达的 → 它上报后，Block Slot 中 0 号 barrier 的 Arrival Mask 终于凑齐，变成了 11111111。

硬件中的比较电路检测到 Arrival Mask == Participation Mask，条件满足，触发全局释放：

d. 所有 warp 继续执行

8 个 warp 重新变成 eligible，各自 partition 的 scheduler 可以从下个 cycle 开始 issue 它们 barrier 之后的指令。

named barrier

GA10x 每个 block slot 里有 16 个独立的 barrier 单元，编号 0–15。

__syncthreads() 默认用 barrier 0。

但你可以通过 cooperative groups 用其他编号的 barrier，让 block 内不同的 warp 子集各自同步而不互相干扰。

每个 barrier 单元都有自己独立的 Arrival Mask（到达位图）和 Participation Mask（参与位图），互不影响。

if 语句

Flynn's Taxonomy, SIMT, and SPMD

SPMD 的意思是，执行同一个程序的时候，处理不同的数据。没有其他更多的约束了。

每一个 CUDA kernel 就是一份 program，launch 时几万个 thread 都跑这同一份 program，但每个 thread 通过 threadIdx、blockIdx 知道自己是谁，从而访问不同的数据。这就是教科书级的 SPMD。

但 GPU 硬件实现混合了多个层次：

所以，在 warp 内部这个角度来看，所有 thread 共享 PC，永远在同一位置，分支靠 mask 模拟（ITS 是例外），可以处理不同的数据，满足 SPMD 定义。

在，一个 partition 上的多个 warp，一个 SM 内多个 partition，多个 SM，这三个角度来看，不同的 thread 有独立 PC，可以走到不同的位置，也可以走到相同的位置，可以处理不同的数据，满足 SPMD 定义。

这就是 NVIDIA 造 SIMT 这个词的原因——编程模型给你 SPMD 的体验（每个 thread 写起来像独立的），硬件用 SIMD 执行（一个 PC 驱动 32 lane）加 MIMD。SIMT 是把 SPMD 编程模型映射到 SIMD/MIMD 硬件上的桥梁。

所以 SPMD 和 SIMD/MIMD 的关系是：SPMD 是程序员看到的抽象，SIMD/MIMD 是硬件干活的方式，SIMT 是 NVIDIA 把两者粘在一起的胶水。