先记住一句话

thread 是编程抽象，warp 是执行/调度基本组，block 是共享 memory 与同步的合作单元，SM 是实际承载 blocks 的硬件。

映射 output 到 threads→组成 warps/blocks→核算 registers/shared memory→估 resident warps→检查 divergence→用 profiler 验证 stalls
Gridall blocks
scheduled onto
SMresident blocks
contains
Warps32 threads
issue
SIMT lanessame instruction
Warp 选择 ready instruction；更多 resident warps 可隐藏等待，但 resources 与 useful work 决定上限。

1. Grid → block → thread

kernel launch = one grid
grid = many independent thread blocks
block = threads that can share shared memory + __syncthreads()
hardware groups consecutive threads into warps (currently 32 threads)

block 必须能独立执行，runtime 才能把 blocks 以任意顺序分配到 SM；不能假设 block 0 一定先于 block 1 完成。

2. SM 里有什么

Streaming Multiprocessor 包含 warp schedulers、execution pipelines、register file、shared memory/L1 与 load/store units；Tensor Cores 是特定 matrix multiply-accumulate pipelines。GPU 有多个 SM，并用大量 resident warps 在一个 warp 等 memory 时切到另一个。

3. SIMT 与 SIMD

SIMT（single instruction, multiple threads）让每个 thread 有自己的 registers、program counter 和地址；warp 中 threads 通常发射相同 instruction。它保留 scalar programming illusion，但不意味着 warp 内 threads 真正完全独立地同时执行所有 divergent paths。

4. Branch divergence

warp 内 threads 对条件选择不同分支时，硬件需执行各 active path 并 mask 另一部分，throughput 下降。若 branch 按 warp/block 粒度一致则几乎无 divergence；短 branch 也可能被 predication 处理。优化要看实际 branch efficiency，不是机械消灭 if。

5. Registers 与 local memory

每 thread registers 最快，但 SM register file 总量有限。kernel 每 thread 用更多 registers，会减少同时 resident warps；超过限制还会 spill 到“local memory”，它物理上通常在 device memory/cache hierarchy，可能很慢。

6. Occupancy

occupancy = active warps per SM / hardware max warps per SM

resident blocks 受 threads/block、registers/thread、shared memory/block 和 architecture limits 共同约束。occupancy 帮助 latency hiding，但 100% 不保证快；compute-heavy kernel 可能在较低 occupancy 已饱和 pipelines，高 occupancy 也可能只让更多 warps 一起 memory stall。

7. Latency hiding 不等于 latency 变小

GPU DRAM load latency 并未消失；scheduler 在 warp A 等待时执行 warp B。若 parallelism 不够、dependency chain 长、所有 warps 同时等 memory，hide 不住。instruction-level parallelism 与 asynchronous copy 也可增加可执行 work。

8. Block size 怎么想

  • 通常取 warp size 的整数倍，避免最后一个 warp 大量空 lanes；
  • 太小：每 block work 少、调度/同步占比大；
  • 太大：register/shared memory 紧张，resident blocks 降低；
  • 二维 tile 要同时匹配 data layout 与 cooperative reuse；
  • 用 occupancy calculator + benchmark，而非永远选 256。

9. Compute capability

不同 GPU architecture 的 SM 资源、Tensor Core types、async copy、shared memory、thread-block cluster 等不同。编译的 SASS/cubin 面向具体 architecture，PTX 可在 driver 侧 JIT；高性能 kernel 常按 compute capability dispatch 多个实现。

10. 四个 SIMT 手算

例 1：grid sizing

N=1,000,000 elements，block=256 threads，需要 ceil(1,000,000/256)=3907 blocks，总 launch threads=1,000,192，最后 192 threads 由 bounds check 退出。

例 2：warp divergence

一个 32-thread warp 中 16 threads 走 20-cycle if、另 16 走 12-cycle else，若两支都执行，warp 约耗 32 cycles 而非 max 20；lane execution efficiency 也会下降。

例 3：register 限制 occupancy

SM 有 65,536 registers，每 thread 用 64 regs，资源最多容纳 65536/64=1024 threads，即 32 warps；若每 thread 32 regs，可到 2048 threads（仍受架构其他上限）。

例 4：latency hiding

Memory latency 400 cycles，warp 每 20 cycles 才再次 ready，理想需约 400/20=20 个独立 warps 轮换填满空档。它隐藏等待，不把单次 load latency 变成 20 cycles。

常见误解：

CUDA thread 数超过 CUDA cores 并不是错误。大量 logical threads 会分批执行，正是用更多 warps 隐藏 latency；真正问题是有没有足够 work、资源和有效 memory access。

自测

1. block 之间能直接用 __syncthreads() 吗？

不能；它只同步同一 block。跨 block 需分 kernel、cooperative groups 的特定机制或 atomics/全局协议。

2. 为什么 register 用多会降低 occupancy？

SM register file 总量固定，每 thread 占用增加后可同时驻留的 threads/warps 变少。

3. divergence 发生在哪个粒度？

同一 warp 内 threads 走不同控制路径时；不同 warps/blocks 走不同 path 不构成同样的 lane masking。

官方资料

NVIDIA Writing SIMT Kernels定义 thread hierarchy、warp execution 与 memory performance；具体资源上限应查目标 GPU 的 compute capability。