GPU架构基础：现代GPU的硬件架构与执行模型
核心原理：从SM到CUDA核心的硬件层次
现代NVIDIA GPU（以Ampere、Hopper架构为例）的硬件设计围绕流式多处理器（Streaming Multiprocessor，SM）展开。SM是GPU最基本的计算单元，一个GPU由多个SM构成，例如A100拥有108个SM，H100拥有132个SM。每个SM内部包含若干CUDA核心（即FP32算术逻辑单元）、Tensor Core、特殊功能单元（SFU）、加载/存储单元（LD/ST）以及共享内存和寄存器堆。
理解SM的组成是内核优化的起点。以A100的SM为例，其内部结构可概括为：
- 4个处理块（Processing Block），每个处理块包含16个FP32 CUDA核心、8个FP64核心、16个INT32核心以及4个Tensor Core。
- 256KB寄存器文件，按线程粒度分配，每个线程最多可使用255个寄存器。
- 164KB共享内存（可配置为L1缓存），支持显式数据复用。
- 64KB L1指令缓存，负责指令分发。
CUDA核心并非独立工作的“小CPU”，而是单指令多线程（SIMT）架构下的执行单元。每个SM中的调度器以warp（32线程）为单位分发指令。当一个warp中的32个线程执行同一指令时，若发生分支发散，则不同路径会被串行执行，导致部分CUDA核心闲置。因此，优化时需尽量避免warp内分支。
线程层次与执行模型
CUDA的线程组织为三层结构：Grid（网格）→ Block（线程块）→ Thread（线程）。内核启动时，Grid被划分为多个Block，每个Block包含若干线程（通常为128～1024的倍数）。硬件调度时，Block被分配到SM上执行，一个SM可同时驻留多个Block（受资源限制）。Block内的线程通过共享内存和同步屏障（__syncthreads()）协作。
执行模型的关键在于warp的调度机制。每个Block内的线程按连续32个为一组形成warp。例如，一个包含256线程的Block会被划分为8个warp。SM中的warp调度器每个时钟周期选择一个就绪的warp发射指令。当某个warp因访问全局内存而等待时，调度器会切换到其他warp执行，以此隐藏内存延迟。这种机制称为延迟隐藏，是GPU高吞吐量的核心来源。
关键参数速查表（以A100为例）
| 参数 | 数值 | 优化影响 |
|---|---|---|
| SM数量 | 108 | 决定并行度上限 |
| 每SM最大线程数 | 2048 | 超过则Block排队 |
| 每SM最大Block数 | 32 | Block太小会浪费调度槽 |
| 每SM最大warp数 | 64 | 影响延迟隐藏能力 |
| 寄存器文件/SM | 256KB | 寄存器溢出会访问本地内存 |
| 共享内存/SM | 164KB | 可配置为L1缓存 |
| L2缓存 | 40MB | 跨SM数据共享 |
内存层次与数据复用策略
GPU内存层次从寄存器到全局内存，延迟差异超过两个数量级。寄存器（<1ns）→ 共享内存（~20ns）→ L2缓存（~200ns）→ 全局内存（~400ns）。优化核心在于最大化数据在高速内存中的复用。
共享内存是程序员可显式控制的高速缓存。典型使用模式包括：
- 分块（Tiling）：将矩阵分块加载到共享内存，减少全局内存访问次数。
- 数据重排：利用共享内存解决bank conflict，提高带宽利用率。
- 合并访问：确保同一warp的线程访问连续内存地址，使每次内存事务传输完整128字节缓存行。
以矩阵乘法为例，朴素实现中每个线程直接读取全局内存，访问次数为O(N³)。使用分块后，每个线程块将子矩阵加载到共享内存，全局内存访问降至O(N³/B)，其中B为块大小。Nsight Compute的Memory Workload分析器可精确显示实际全局内存吞吐量与理论峰值的比值，帮助判断是否达到合并访问要求。
案例：通过Nsight Compute分析SM占用率
假设我们有一个内核，启动配置为grid(256), block(256)，每个线程使用40个寄存器。在A100上，每SM可驻留的线程数受限于：
- 寄存器限制：65536个寄存器/SM ÷ 40 = 1638线程
- 线程数限制：2048线程/SM
- Block数限制：32 Block/SM
实际最大驻留线程数为1638（取三者最小值），SM占用率为1638/2048 = 80%。Nsight Compute的Launch Statistics页面会显示Theoretical Occupancy和Achieved Occupancy。若实测占用率远低于理论值，通常是因为Block调度不均匀或warp停滞。
# 使用ncu查看内核占用率
ncu --section LaunchStats --section Occupancy ./my_kernel
# 输出关键指标示例
# Achieved Occupancy: 0.78
# Theoretical Occupancy: 0.80
# Block Limit: 32
# Registers Per Thread: 40
# Shared Memory Per Block: 0 bytes若占用率低于预期，可采取以下措施：
- 减少每线程寄存器使用（通过
__launch_bounds__或重构代码） - 调整Block大小以匹配SM的warp调度粒度
- 检查是否存在共享内存或寄存器导致的驻留限制
误区与复盘：常见优化陷阱
误区1：盲目追求100%占用率
高占用率并不总是带来高性能。当内核存在大量共享内存访问或同步操作时，过高的占用率会导致warp间竞争资源，增加bank conflict和同步延迟。Nsight Compute的Warp State Statistics可显示warp停滞原因——若大量warp因Long Scoreboard（等待内存）而停滞，增加占用率才有意义；若因Barrier或MIO Throttle停滞，则应减少Block规模或优化同步逻辑。
误区2：忽略warp发散的影响
在SIMT架构下，若warp内线程执行不同分支，所有路径都会被串行执行。例如，一个warp中16个线程进入if分支，16个进入else分支，则两条路径各执行一次，总耗时翻倍。Nsight Compute的Control Flow部分会显示分支效率。优化方法包括：将条件判断移到warp级别、使用位运算替代分支、或重新组织数据使同一warp的线程走相同路径。
误区3：共享内存使用不当
共享内存虽然快，但容量有限且存在bank冲突。当多个线程同时访问同一bank的不同地址时，访问会被串行化。例如，shared[tid * 32]会导致所有线程访问同一bank，产生32路冲突。正确做法是使用shared[tid]或添加填充（padding）来避免冲突。Nsight Compute的Shared Memory分析器会报告bank conflict次数，帮助定位问题。
复盘建议：每次内核优化后，使用Nsight Compute的Memory Workload Analysis和Compute Workload Analysis对比优化前后的指标变化。重点关注：全局内存吞吐量、共享内存bank conflict、warp停滞周期、指令混合比。只有当这些指标与理论瓶颈匹配时，优化才算真正生效。例如，若内核是内存密集型，优化后全局内存吞吐应接近峰值；若为计算密集型，则FP32利用率应显著提升。