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

要写出高性能的CUDA内核，必须先建立一张清晰的“硬件地图”。GPU并不是一堆同质的核心简单堆叠，而是一个分层、分级的执行与存储体系。理解SM、SP、Warp、线程块、网格这些执行单元，以及寄存器、共享内存、全局内存、常量内存、纹理内存、L1/L2缓存这些存储层级，是后续做占用率调优、内存合并访问、分支发散治理的前提。

执行层级：从网格到线程

CUDA程序的执行组织是四层结构。最外层是网格（Grid），对应一次内核启动；网格由若干线程块（Block）组成；每个线程块由若干线程（Thread）组成。硬件侧，一个GPU由多个流多处理器（SM，Streaming Multiprocessor）构成，每个SM内部包含若干流处理器（SP，Streaming Processor，也称CUDA Core）、寄存器文件、共享内存/L1缓存、Warp调度器和加载/存储单元。

关键映射关系是：线程块被整体调度到一个SM上执行，块内线程不再跨SM迁移；而Warp是硬件调度和执行的最小单位，32个连续线程被编成一个Warp。也就是说，线程块是“资源分配与驻留单位”，Warp是“指令发射单位”。一个块即使只有1个线程，也会占用一个Warp的槽位，其余31个通道被浪费。

举例：启动 kernel<<<dim3(64,4), dim3(128)>>>()，即网格有256个块，每块128线程。每块128线程 = 4个Warp。若SM最多驻留16个块、64个Warp，则理论上该SM可同时容纳16个块（64个Warp），但受寄存器、共享内存限制，实际驻留数往往更少。

存储层级：速度、容量与作用域

存储体系按“越靠近SP越快、越小、越贵”排列：

• 寄存器：每个线程私有，速度最快，容量最小。内核中局部变量默认放寄存器，寄存器用量直接决定每SM可驻留的线程数。
• 共享内存（Shared Memory）：块内共享，片上SRAM，延迟远低于全局内存。常用于块内数据复用、矩阵分块、归约。需注意bank冲突。
• L1缓存：与共享内存共用同一块物理SRAM（可配置划分），缓存全局内存访问。
• L2缓存：全GPU共享，容量较大，是全局内存访问的统一汇聚点。
• 全局内存：容量最大、延迟最高（数百周期），所有线程可访问，需靠合并访问和缓存提升有效带宽。
• 常量内存：只读，有专用常量缓存，适合所有线程读取同一地址的场景（如系数、滤波器权重），广播时极快。
• 纹理内存：只读，经纹理缓存，适合具有空间局部性的二维访问，自带插值和边界处理。

下面这张图把执行层级与存储层级放在同一张“硬件地图”上，帮助建立整体印象：

内存合并访问：带宽的第一道闸门

全局内存访问以“内存事务”为单位。当一个Warp的32个线程访问的地址落在同一段连续区间内，硬件可将其合并为少量事务；若地址分散，则退化为多次事务，有效带宽骤降。

以 float 数组为例，Warp中线程 t 访问 data[base + t]，32个线程覆盖128字节连续区间，通常合并为1次128字节事务，效率最高。若改为 data[base + t*32]，地址间隔128字节，32个线程分散在32段不同区间，可能产生32次事务，带宽利用率降到约1/32。因此写内核时，应让相邻线程访问相邻地址，并尽量让起始地址对齐到128字节边界。

分支发散：Warp内的串行化代价

Warp内32个线程共享一个指令流。遇到 if/else 时，如果线程分成两组走不同分支，硬件会先执行一组、屏蔽另一组，再执行另一组，总耗时接近两条路径之和。这就是分支发散。

例如 if (tid % 2 == 0) {...} else {...}，Warp内奇偶线程各半，两条路径都要执行，效率约减半。若条件只与Warp编号相关（如 if (warpId < 4)），则整个Warp走同一路径，无发散。优化思路：把分支条件尽量提升到Warp粒度，或用无分支写法（如 select、掩码运算）替代。

占用率：延迟隐藏的调节旋钮

GPU靠大量Warp切换来隐藏内存与指令延迟。占用率衡量每SM实际驻留的Warp数占最大值的比例。它由三方面约束：

• 寄存器：每线程寄存器数 × 线程数 ≤ SM寄存器文件大小。寄存器越多，可驻留线程越少。
• 共享内存：每块共享内存 × 驻留块数 ≤ SM共享内存容量。
• 块与Warp上限：SM对驻留块数、Warp数有硬上限。

可用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 估算，或编译时加 --ptxas-options=-v 查看寄存器与共享内存用量。注意：占用率不是越高越好。若内核是计算密集型且寄存器复用充分，适度降低占用率换取更多寄存器反而更快；若内核是访存密集型，则需要较高占用率来掩盖延迟。

一个可对照的示例

假设在示例GPU上，SM寄存器文件为65536个寄存器，共享内存为100KB，最大驻留64个Warp。某内核每线程用40个寄存器、每块用8KB共享内存、块大小256线程（8个Warp）。则：

• 寄存器约束：65536 / (40 × 256) ≈ 6.4 → 最多6个块；
• 共享内存约束：100KB / 8KB = 12个块；
• Warp上限：64 / 8 = 8个块。

取最小值，每SM驻留6个块 = 48个Warp，占用率约 48/64 = 75%。若把每线程寄存器压到32个，则寄存器约束变为 65536/(32×256)=8个块，占用率可达100%。这说明寄存器用量是占用率最敏感的杠杆之一。

把这些层级关系、映射规则和约束条件记牢，后续做内核优化时，就能快速判断瓶颈出在执行组织、访存模式还是资源占用上。