异构系统

CUDA 编程模型是基于 CPU + GPU 的异构计算系统。与 CPU 直连的内存被称为主机内存 host memory，与 GPU 直连的内存称为设备内存 device memory。

CUDA 程序会在 GPU 上执行部分代码，但整个程序的执行仍然始于 CPU。在 CPU 上运行的代码称为 host 代码，在 GPU 上运行的代码称为 device 代码。host 代码可以使用 CUDA API 在主机内存和设备内存之间传输数据，启动在 GPU 上的代码执行，并等待数据传输或 GPU 代码的完成。

CPU 和 GPU 可以同时执行代码，以获得最佳性能。

GPU 硬件结构

这是一张 NVIDIA GH100 GPU 的架构图，来源于《NVIDIA H100 白皮书》中的 Figure 6。从上而下的，依次查看 GPU 中的硬件构成。

GPU 硬件架构总览

SM 架构

这是一个 SM 内部的硬件架构示意图，来源于《NVIDIA H100 白皮书》中的 Figure 7。从上而下的，依次查看 SM 中的硬件构成。

此外，我们常说的片上、片下，指的是整块 GPU 硅片。「片上」是 on-chip 的中文译法，指在 GPU 芯片内部，与之相对的是 off-chip「片外」，分界线是显存 HBM3。

• 片上 on-chip：L1/L0 指令缓存、共享内存、L1 数据缓存、寄存器、L2 缓存，都刻在同一块硅片上
• 片外 off-chip：HBM3 显存，是独立的存储芯片，通过封装基板与 GPU 相连

访存延迟

一次显存访问的时间，够执行几百条算术指令。如果一个 warp 发出访存请求后就一直等待，那这几百个周期里执行单元完全闲置。

CPU 通过完善的控制逻辑来应对访存延迟，如大容量缓存、分支预测、乱序执行、预取等，减少访存次数并避免停顿。而 GPU 则通过海量 warp 轮转来掩盖访存延迟，当某个 warp 因访存阻塞时，SM 的 Warp Scheduler 立刻转向下一个就绪的 warp。GPU 掩盖访存延时依赖于 warp 切换的零开销。

CPU 线程切换要保存或恢复寄存器上下文，代价上千周期。而 GPU 上，每个 warp 的寄存器是一直独占的、不搬走的。这正是每个 SM 中 64 K 寄存器的用途，64 个 warp 的上下文全部常驻在寄存器堆里。切换只是调度器换一个索引，不涉及任何数据搬运，因此不花周期。

SM 的 warp 约束

寄存器容量约束

为了 GPU 线程切换的零开销，每个线程的上下文必须常驻在寄存器中。而数据一旦常驻，寄存器的容量就成为约束。

每个 SM 上可驻留的线程数 = 每个 SM 的寄存器总个数 / 每线程占用寄存器个数 = 每个 SM 的寄存器总容量 / 每线程占用寄存器容量

以 H100 举例，每个 SM 的寄存器个数为 16384 * 4 = 65536 个，寄存器总容量为 64KB * 4 = 256 KB。

如果每线程占用 32 个寄存器，每线程占用寄存器容量 32 * 32 b = 0.125 KB，则每个 SM 上可驻留的线程数 = 64K / 32 = 256 KB / 0.125 KB = 2048 个，换算成 warp 为 2048 / 32 = 64 个。

共享内存容量约束

SM 细节图中的「256 KB L1 Data Cache / Shared Memory」（序号 ⑫）部分，是共享内存（程序显式分配）与 L1 数据缓存（硬件自动管理）所在地。共享内存是同一 block 内线程共享数据的物理载体。

共享内存的生命周期等于 block 的生命周期，从 block 被调度到 SM，直到 block 执行结束，这块内存一直分配给该 block，期间不能被换出或让给其他 block。这与寄存器同理，为保证 warp 切换零开销，驻留对象的状态必须常驻片上，因此共享内存的容量也成为驻留数量的约束。

以 H100 举例，SRAM 总共 256KB，其中共享内存最多可配置到 228KB，划分比例可使用 cudaFuncSetAttribute 调整。

需注意 CUDA 默认每个 block 最多能占用 48KB 共享内存，超过 48KB 必须使用动态分配（extern __shared__ + 启动参数）并通过 cudaFuncSetAttribute 显式 opt-in，否则 kernel 启动失败。

假设声明每个 block 占用的共享内存大小为 114KB，每个 SM 上最多可驻留 228KB / 114KB = 2 个 block。实际驻留的线程数还取决于 block 大小，假设每个 block 含 256 个线程，即最多可驻留 512 个线程，换算成 warp 为 512 / 32 = 16 个。warp 占有率 = 实际驻留的 warp 数 / 硬件上限 = 16 / 64 = 25%。

block 槽位约束

每个 SM 要同时管理多个驻留 block，需要给每个 block 分配一套管理资源，如 block 索引、同步屏障计数器、共享内存基址寄存器等。这些管理槽位的数量在芯片设计时就已经确定。block 槽位就是一个 SM 能同时「记住」多少个 block 的位置数。与寄存器、共享内存不同，它不是「容量 / 每线程占用量」算出来的，而是硬件的记账位数上限。

以 H100 为例，每个 SM 最大的 block 槽位是 32 个，即同一时刻最多有 32 个 block 驻留在同一个 SM 上。

总结

CUDA 编程模型

SIMT

SIMT：Single Instruction, Multiple Threads（单指令多线程），NVIDIA 给 GPU 执行模型起的名字。

SIMD：Single Instruction, Multiple Data（单指令多数据），CPU 的向量指令，一条指令同时对多个数据做同样的运算。

SIMT 的核心在于，一条指令同时驱动一个 warp 的 32 个线程执行，每个线程处理不同的数据。

SIMT 的三个特点

1. 指令是共享的，数据是各自的
2. 写代码时是单线程视角
3. 线程之间靠下标区分

锁步

锁步（lock-step），SIMT 的执行方式。而 warp 是 SIMT 的作用单位。

由于一个 warp 只有一套取指和译码逻辑，因此需要一个 warp 的 32 个线程像被锁在一起一样，同进同退，同一时刻执行同一条指令。

分支发散

分支发散，SIMT 的代价。

由于一个 warp 的 32 个线程只有一套取指和调度逻辑，同一时刻只能执行一条指令。所以当同一 warp 内的 32 个线程要走不同分支时，硬件没法做到让它们真正同时跑，只能分多趟。

if (threadIdx.x % 2 == 0) {
    A();     // 偶数号线程走这里
} else {
    B();     // 奇数号线程走这里
}

例如上面的 cuda 代码，让一个 warp 里 16 个线程走 A、16 个走 B。

硬件的做法是，第一趟：执行 A。16 个偶数线程正常计算，16 个奇数线程被屏蔽（lane 空转，什么也不做）。第二趟：执行 B。16 个偶数线程被屏蔽，16 个奇数线程正常计算。总耗时 = A 的时间 + B 的时间，而不是 max(A, B)，即便每一趟都有一半线程在闲着。

分支发散只发生在 warp 内部，不同 warp 走不同分支完全没有代价，因为它们本来就是独立调度的。这也是优化分支发散的核心思路，让分支条件以 warp 为粒度对齐，而不是以线程为粒度。常见写法是把条件改成 (idx / warpSize) % 2 == 0。

Thread、Lane、CUDA Core 的区别

CUDA 软件层级与 GPU 硬件架构的对应关系

附件

• CUDA C++ Programming Guide
• 《NVIDIA H100 Tensor Core GPU Architecture Whitepaper》：本文 GPU 硬件结构与 SM 架构两节的 Figure 6、Figure 7 均出自该白皮书