CUDA 编程模型基础
一句话定位：理解 GPU 的两件事——线程怎么被组织和调度、数据放在哪一层存储。用大数据的调度与分层存储直觉去套，学习成本最低。
1. 线程层次：Grid → Block → Thread
- Grid：一次 Kernel 启动的全部线程，由多个 Block 组成。
- Block：线程块，块内线程可通过共享内存通信与
__syncthreads()同步；一个 Block 必须整体调度到同一个 SM 上。 - Thread：最小编程单位，通过
blockIdx * blockDim + threadIdx计算全局索引。
关键区分（面试易错点）：
- Warp（32 线程）是真实的调度单位。硬件以 warp 为粒度取指与执行（SIMT），warp 内 32 线程执行同一条指令。因此若 warp 内线程走了不同分支，会产生warp divergence（分支被串行执行），性能下降。Block 大小通常取 32 的倍数就是这个原因。
- SM（Streaming Multiprocessor）是执行硬件：包含 CUDA Core、寄存器文件、共享内存、warp 调度器。多个 Block 可并发驻留同一 SM（占用率 occupancy 由寄存器/共享内存用量决定）。
2. 显存层次（速度差数量级）
从快到慢：
| 层级 | 作用域 | 相对速度 |
|---|---|---|
| 寄存器 (Register) | 单线程私有 | 最快（~1 周期） |
| 共享内存 (Shared Memory) | Block 内共享，片上 SRAM | 很快（~几十周期） |
| L2 Cache | 全 GPU 共享 | 中等 |
| 全局内存 (Global / HBM) | 全 GPU + Host 可见 | 最慢（数百周期） |
核心结论：优化 GPU 程序的主线是”把数据尽量留在高层级”——用共享内存复用数据、保证全局内存合并访问（coalesced），减少对 HBM 的读写次数。这正是 FlashAttention（3.7）的核心思想。
3. 可迁移类比（我的表达桥梁）
- Warp 调度 ≈ Spark Task 调度：都是把大任务切成固定粒度的执行单元批量调度；warp divergence ≈ Task 内分支导致的负载不均。
- 显存层次 ≈ 内存/磁盘分层存储：寄存器/共享内存 ≈ 内存缓存，全局内存(HBM) ≈ 磁盘/远端存储；”减少 HBM 读写” ≈ “减少 Shuffle 落盘与网络传输”，都是把 IO 成本而非算力当作首要瓶颈。
参考：《CUDA C++ Programming Guide》