2. GPU硬件架构与CUDA执行模型

要写出真正跑得快的CUDA程序，必须先弄清楚代码里的 threadIdx、blockIdx 到底映射到硬件的哪一层，以及一个线程从发射到完成要经过哪些资源。本节把GPU的物理层级与CUDA的编程层级对齐来看，并解释占用率、延迟隐藏这些工程概念背后的硬件原因。

从网格到SM：两级映射

CUDA的编程模型是三层结构：线程（thread）→ 线程块（block）→ 网格（grid）。硬件侧对应的执行单元是：线程 → warp → 线程块 → SM（流多处理器）。一个kernel启动后，grid被拆成若干block，block被调度器分配到各个SM上执行；一个SM可以同时驻留多个block，但一个block一旦被分配，其所有线程都只在这一个SM上运行，不会跨SM拆分。这一点决定了共享内存和 __syncthreads() 的语义边界——它们的作用域就是block。

线程在硬件上并不是逐个调度的，而是以warp为单位，一个warp固定包含32个线程。block内的线程按 threadIdx 线性编号，每32个连续线程组成一个warp。因此block大小最好取32的整数倍，否则最后一个warp会有线程被"浪费"（inactive lane），既降低效率又浪费寄存器配额。

SIMT执行：一条指令，32个线程

SIMT（单指令多线程）是GPU的核心执行方式：一个warp内的32个线程在同一时刻执行同一条指令，但各自操作自己的寄存器数据。这带来一个关键约束——分支发散。如果warp内线程走了不同的 if 分支，硬件只能串行执行各分支路径，未走该路径的线程被屏蔽。例如：

if (threadIdx.x % 2 == 0) {
    a = a + 1;      // 偶数线程执行，奇数线程等待
} else {
    a = a * 2;      // 奇数线程执行，偶数线程等待
}

这段代码在warp内会执行两次，实际吞吐减半。工程上的应对是尽量让分支条件在warp粒度上一致（如按 threadIdx.x / 32 判断），或改用无分支的算术/查表写法。

寄存器与共享内存：两级片上存储

每个线程拥有私有的寄存器，访问延迟约1个周期，是最快的存储。寄存器数量有限（如每个SM 65536个32位寄存器），一个线程用得多，能同时驻留的线程就少。编译器会尽量把局部变量放进寄存器，但数组若被动态索引，就可能被"溢出"到本地内存（实际在显存中），延迟骤增。

共享内存是block内所有线程共享的片上存储，延迟约20~30周期，远低于全局内存的400~800周期。它常用于：块内数据交换、矩阵分块（tiling）、归约等。共享内存按bank划分（通常32个bank，每bank 4字节宽），若同一warp内多个线程访问同一bank的不同地址，就会发生bank冲突，访问被串行化。经典规避手段是给二维数组的列维加padding，例如把 __shared__ float tile[32][32] 改为 [32][33]。

下表对比三类存储的关键参数（数值为典型量级，具体依架构而定）：

占用率与延迟隐藏

GPU靠线程级并行来隐藏延迟：当一个warp在等待全局内存返回时，SM的warp调度器立刻切换到另一个就绪warp执行。因此，SM上同时驻留的warp越多，越容易"填满"等待空隙，这就是延迟隐藏。

占用率（occupancy）定义为：SM上实际驻留的warp数 ÷ 该SM支持的最大warp数。它受三个资源约束——寄存器、共享内存、block数量上限。以某架构SM为例（示例参数）：最大2048线程、64个warp、65536个寄存器、共享内存可配。若某kernel每线程用40个寄存器，则寄存器能支持的线程数为 65536/40 ≈ 1638，向下取整到warp粒度约51个warp，占用率约 51/64 ≈ 80%。若把寄存器压到32个，则可支持64个warp，达到100%。

但占用率不是越高越好。有时用更多寄存器换取更少的指令、更少的重复访存，反而更快。工程判断顺序应是：先保证访存合并、减少发散，再观察占用率是否成为瓶颈。可用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 估算，或通过 --ptxas-options=-v 查看寄存器用量。

工程含义：把模型落到调优动作上

理解这套模型后，调优就有了明确抓手。第一，block大小取128或256是常见起点，既能凑满warp，又给SM留出驻留多个block的余地。第二，合并访存：让同一warp的32个线程访问连续的全局内存地址，硬件可合并为少数几次事务；若地址跨步，事务数成倍增加。第三，减少发散：把分支条件提到warp粒度。第四，善用共享内存做tiling：矩阵乘法中每个block把一块A、B载入共享内存，块内线程反复复用，把全局访存次数从O(N³)降到O(N²)。第五，用占用率做诊断而非目标：当性能不随占用率提升而改善时，瓶颈往往在访存模式或指令依赖，而非并行度不足。

下面给出一个可直接套用的检查清单，用于快速定位kernel的硬件层瓶颈：

• 用 nvcc -Xptxas -v 查看每线程寄存器数与共享内存用量，判断是否限制了驻留warp数。
• 用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 计算理论占用率，与实测对比。
• 检查全局访存是否连续：同一warp内线程的地址应相差1个元素。
• 检查共享内存是否bank冲突：二维数组列维加1做padding。
• 检查分支是否在warp内发散：把条件改为按 threadIdx.x/32 判断。
• 用Nsight Compute查看 sm__warps_active 与 l1tex__t_sectors，确认是占用率问题还是访存问题。

把编程层级与硬件层级对齐，是后续所有优化工作的基础。记住三个数字：warp=32、block不跨SM、全局内存延迟是共享内存的十几倍。围绕它们设计数据布局与并行划分，性能提升往往立竿见影。