1. GPU 整体鸟瞰

GPU 由多个 Streaming Multiprocessor (SM) 组成，SM 通过 L2 Cache 和 DRAM 交互。点击任意 SM 查看提示。

关键组件

• SM：GPU 的计算核心，同时驻留多个 block/warp。
• DRAM：大容量全局显存，速度相对慢。
• L2 Cache：所有 SM 共享，缓冲 DRAM 访问。
• PCIe / NVLink：与 CPU / 其他 GPU 通信。

数据流动

Thread ↔ Register ↔ Shared Memory/L1 ↔ L2 ↔ DRAM。越靠近 thread，越快、越小。

2. SM 内部结构

一个 SM 就像一个小型工厂：Warp Scheduler 是调度员，CUDA Cores 是工人，Register File 和 Shared Memory 是车间里的快速缓存。

组件说明

• Warp Scheduler：每周期从 eligible warps 中选一个，把指令发给 Dispatch Unit。
• Dispatch Unit：把指令派发到具体的执行单元。
• CUDA Core：执行单个 thread 的标量指令。一个 SM 通常有 64~128 个 FP32 CUDA Cores。
• Register File：给所有活跃 thread 分配寄存器，容量大、延迟极低。
• Shared Memory / L1：同 block 内 thread 共享，可手动管理。
• Load/Store Unit：负责 global/local memory 访问。
• SFU / Tensor Core：处理特殊函数和矩阵乘/卷积加速。

3. 执行层级：Grid → Block → Warp → Thread

关键数字

• Warp Size：固定 32 个 threads。
• Block Size：通常是 32 的倍数（64, 128, 256, 512, 1024）。
• 一个 Block 内的 threads 可以共享 Shared Memory 并同步（__syncthreads）。
• 不同 Block 之间不能同步，也不能直接共享内存。

索引公式

1D：global_id = blockIdx.x * blockDim.x + threadIdx.x

2D：row = blockIdx.y * blockDim.y + threadIdx.y; col = blockIdx.x * blockDim.x + threadIdx.x

4. Warp Scheduler：隐藏延迟的秘密

当某个 warp 等待内存数据时，Warp Scheduler 会立刻切换到另一个 ready warp，让 CUDA Cores 几乎不空闲。

调度过程

1. Kernel launch 后，grid 被切成多个 block。
2. Block 被分配到空闲的 SM 上。
3. 每个 block 内部的 threads 被进一步切成 32-thread 的 warp。
4. Warp Scheduler 每周期检查 eligible warp pool，选一个 ready 的 warp。
5. 如果选中的 warp 需要访问内存，把它移到 waiting pool，同时调度另一个 warp。
6. 内存数据到达后，waiting warp 回到 eligible pool。

为什么这叫 latency hiding？ 内存访问可能需要几百个时钟周期。Warp Scheduler 通过快速切换 warp，让 CUDA Cores 总有事情做。

5. SIMT、Warp Divergence 与 Memory Coalescing

如何避免 divergence？

• 让同一个 warp 内的 threads 走相同分支。
• 如果必须分支，尽量让分支边界与 warp 边界对齐。
• 利用 Triton/CUDA 的 ballot/shuffle 等 warp 级原语。

6. Occupancy 计算器

Occupancy = 每个 SM 上同时活跃的 warp 数 / SM 支持的最大 warp 数。它决定了 latency hiding 的能力。

常见瓶颈

• threads：block size 太大，每个 SM 放不下足够多的 block。
• registers：每个 thread 用太多寄存器。
• shared memory：每个 block 用太多 shared memory。
• warps：同时需要的 warp 数超过 SM 上限（较少见）。