CPU 与 GPU 的取舍

CPU 用大量晶体管优化少量线程的低延迟：复杂控制、乱序执行、大缓存、强分支预测。GPU 把更多面积交给吞吐：大量算术单元、宽内存系统和硬件线程调度，用许多并发 warp 隐藏单个操作的长延迟。

GPU 不是“每个线程都很快”；它依靠足够多的线程 + 快速切换就绪 warp维持流水线。

从 GPU 到线程的层次

GPU Device
├─ GPC / 更高层分区（具体组织因架构而异）
│  └─ SM 0
│     ├─ warp schedulers / dispatch units
│     ├─ FP/INT CUDA Cores
│     ├─ Tensor Cores
│     ├─ Load/Store 与 Special Function Units
│     ├─ register file
│     └─ L1 / shared memory
├─ SM 1 ... SM N
├─ L2 cache（所有 SM 共享）
├─ memory controllers
└─ GDDR/HBM device memory

这张图是编程心智模型，不是所有芯片的精确版图。每代 GPU 的 SM 内部分区、执行单元数量、指令吞吐、L1/shared 组合方式都可能变化；写可移植代码时依赖 Programming Guide 的抽象，调优时再看对应架构文档和实测指标。

SIMT 与 warp

CUDA 对程序员暴露独立线程，硬件把线程按 warp 分组执行。当前 CUDA 支持的架构中 warp size 是 32，但代码仍应使用 warpSize 或库抽象表达意图。

一个 warp 的 32 个 lane
lane:  0  1  2  3  4  5 ... 30 31
       └──────── 同一条 warp 指令 ────────┘

if (lane < 16) A(); else B();
时间： [A: lane 0..15 active] [B: lane 16..31 active]
       inactive lane 不做有效工作，但路径仍需被执行

这就是 SIMT（Single Instruction, Multiple Threads）：线程拥有独立寄存器状态、地址和控制流，但调度和发射以 warp 为核心。与传统 SIMD 的区别在于程序仍写“线程”，硬件处理活跃掩码与路径。

分支发散并不等于出现 if

关键是同一 warp 内的线程是否走不同路径：

// 通常不发散：不同 warp 走不同分支。
if (threadIdx.x / warpSize == 0) { ... }

// 高概率发散：相邻 lane 交替走分支。
if ((threadIdx.x & 1) == 0) { ... } else { ... }

短分支可能被谓词化；真实代价取决于两侧指令、活跃线程比例、编译器与架构。不要仅凭源码猜，查看 SASS 和 Nsight Compute 的 branch/warp efficiency 指标。

SM 如何隐藏延迟

一个 warp 发出 global load 后可能等待很长时间。SM 不保存/恢复操作系统线程那样的上下文，而是让许多 warp 的寄存器状态同时驻留，迅速选择另一个就绪 warp：

cycle 0: Warp A 发出 load ────────────────┐
cycle 1: Warp B 做 FMA                    │ A 等数据
cycle 2: Warp C 做 integer op             │
cycle 3: Warp D 发出 load                 │
...                                       │
cycle k: Warp A 数据返回，继续计算 ◄──────┘

能够驻留多少 warp 受以下资源共同限制：

• 每 SM 最大 warp/thread/block 数；
• block 大小；
• 每线程寄存器数；
• 每 block 静态与动态 shared memory；
• cluster 等架构专属约束。

Occupancy 定义为：

$$\text{Occupancy}=\frac{\text{active warps per SM}}{\text{maximum warps per SM}}$$

它衡量“有多少候选 warp 可供调度”，不是“执行单元利用率”。低 occupancy 常会削弱隐藏延迟的能力；高 occupancy 仍可能因为所有 warp 同时等待内存而很慢。

Block 为什么不能随意跨 SM

一个 block 的所有线程被安排在同一个 SM 上直到完成，原因是它们需要：

• 低成本共享 shared memory；
• 使用 block barrier 同步；
• 共同占用固定的寄存器和 shared memory 配额。

多个 block 可以并发驻留在同一 SM；一个 block 不会执行到一半迁移到另一个 SM。不同 block 的执行次序默认未定义，普通 kernel 内不能假设 “block 0 一定先于 block 1”。

Hopper（Compute Capability 9.0）开始提供 Thread Block Cluster 与 Distributed Shared Memory，但这是显式、受能力与 launch 配置约束的高级机制，不改变“普通 block 默认独立”的基础模型。

计算单元地图

“CUDA Core 数量”不能单独预测实际性能。还要考虑频率、每周期吞吐、数据类型、内存带宽、SM 数、功耗状态、指令混合、占用率和工作负载。

Compute Capability 是功能合同

Compute Capability（CC）定义设备支持的能力与资源边界。当前官方表覆盖 7.x、8.x、9.0、10.x、11.0、12.x；一些关键演进：

例如 sm_90、sm_90a 不是同一个兼容承诺。带 a 的目标允许使用架构专属特性，只保证在精确目标上运行；带 f 的目标表达家族内共享特性；无后缀是更广的 baseline feature set。发行包应同时考虑功能需求和兼容范围。

硬件上限不要写死。查询：

nvidia-smi --query-gpu=name,compute_cap,memory.total --format=csv

cudaDeviceProp p{};
cudaGetDeviceProperties(&p, 0);
std::cout << p.name << " cc=" << p.major << '.' << p.minor
          << " SM=" << p.multiProcessorCount
          << " warp=" << p.warpSize
          << " maxThreads/block=" << p.maxThreadsPerBlock
          << " shared/block=" << p.sharedMemPerBlock << '\n';

架构演进应怎样影响代码

保持三层结构：

算法语义层：结果是什么？          CPU reference / 测试
可移植 CUDA 层：大多数 GPU 都能跑  baseline kernel / CUDA-X library
架构特化层：特定 GPU 更快           async copy / TMA / cluster / 专属 Tensor Core

运行时根据 cudaGetDeviceProperties 或编译期根据 __CUDA_ARCH__ 选择路径。不要把整个算法绑定到某代 GPU 的固定 SM 数、固定 shared memory 容量或固定 Tensor Core tile。

本章自测

• 为什么 GPU 在单线程延迟上通常不如 CPU，却在大规模规则计算上更快？
• warp 发散为什么只关心同一 warp 内的路径差异？
• 高 occupancy 为什么不保证高性能？
• block 使用过多寄存器或 shared memory 时，会怎样影响 SM 上的并发 block 数？
• sm_90 与 sm_90a 的兼容承诺有什么不同？

官方参考

• CUDA Programming Guide: Compute Capabilities
• CUDA Programming Guide: SIMT Kernels
• Hopper Tuning Guide
• Blackwell Tuning Guide