CUDA 线程层次映射 — 软件配置 → 逻辑层次 → 物理硬件
1. 软件配置
kernel<<<gridDim, blockDim>>>()
gridDim = (N + blockDim - 1) / blockDim
总共启动多少个 block — 可以任意大, GPU 自动分批调度
blockDim = 256 (通常取 128 / 256 / 512)
每个 block 多少线程 — 必须是 32 的倍数, 上限 1024
2. 逻辑层次
Grid (网格)
— 全部 block 的集合, 分布在 GPU 所有 SM 上。不同 block 之间不能同步, 不能访问对方 Shared Memory。
Block 0
256 threads, 1 个 SM
Block 1
256 threads, 1 个 SM
Block 2
256 threads, 1 个 SM
... 共 gridDim 个 block
Block 0 内部展开:
Warp 0
Thread 0–31
Warp 1
Thread 32–63
Warp 2
Thread 64–95
... ×8 warps = 256 threads
blockDim = 256 → 256 threads = 8 warps。每 32 线程自动分组为一个 warp, 由 Warp Scheduler 锁步执行 (SIMT)。
3. 物理硬件
GA102 GPU (RTX 3090) — 82 SM
SM0
SM1
SM2
... ×82 SM
SM 内部结构
Warp Scheduler ×4
CUDA Core ×128
Tensor Core ×4
Register / L1 / Shared
对应关系
Grid = 整个 GPU 芯片
gridDim 个 block 分布到所有 SM。自动分批 (wave) — block 0~821 先跑, 然后 822~, ...
Block = 1 个 SM
一个 block 固定在一个 SM 上执行，不跨 SM。block 内线程通过 Shared Memory 通信并同步。
Warp = Warp Scheduler 的调度单位
每 32 线程 = 1 warp (硬件自动分组)。
SM 内 4 个 Warp Scheduler，每周期各发射 1 个 warp → 最多 128 threads/cycle。
一个 warp 等内存 (stall ~600 cycles) 时，Scheduler 零开销切到另一个就绪 warp — GPU 靠此隐藏延迟。
Thread = 1 个 CUDA Core 上的 1 条指令
RTX 3090 (1536 threads/SM): blockDim=256 → 6 blocks/SM → 100% | blockDim=1024 → 1 block/SM → 67%
H100 (2048 threads/SM): blockDim=256 → 8 blocks/SM → 100% | blockDim=1024 → 2 blocks/SM → 100% | blockDim 和 GPU 型号相关, 不是绝对值