GPU架构基础：从硬件结构到并行计算模型
核心原理：GPU为何“宽”而非“深”
CPU与GPU的设计哲学差异，决定了它们各自适用的计算场景。CPU仅有少数几个（通常4~16个）高性能核心，每个核心拥有复杂的乱序执行逻辑、大容量缓存和分支预测器，擅长处理延迟敏感、依赖密集的串行任务。GPU则集成了数千个轻量级计算核心，它们共享简单的调度单元，以极高的线程吞吐量掩盖访存延迟，专为数据并行（Data Parallelism）而设计。理解这一本质区别，是后续所有性能分析工作的前提。
关键认知：GPU的“快”并非来自单核频率，而是来自“人多力量大”——当数百万个线程同时执行相同指令时，即便每个线程运行缓慢，整体吞吐量依然惊人。因此，GPU性能优化的核心目标，是让尽可能多的线程保持忙碌，而非减少单个线程的执行时间。
GPU硬件架构：SM、CUDA核心与内存层次
现代NVIDIA GPU（如Ampere、Hopper架构）由多个图形处理簇（GPC）构成，每个GPC内含若干流式多处理器（Streaming Multiprocessor，SM）。SM是GPU最基本的计算单元，相当于一个“小型的多核CPU”，但内部结构截然不同。一个典型的SM包含以下关键组件：
- CUDA核心（FP32单元）：执行单精度浮点运算的算术逻辑单元，每个SM包含64~128个不等。它们被组织为多个处理块（Processing Block），共享指令调度逻辑。
- 特殊功能单元（SFU）：执行超越函数（如sin、cos、log、倒数等），吞吐量通常低于CUDA核心。
- 加载/存储单元（LD/ST）：负责在SM与各级缓存/显存之间搬运数据，其数量直接影响访存带宽。
- 共享内存（Shared Memory）：位于SM内部的可编程高速缓存，容量通常为48KB~228KB（可配置），供同一线程块内的线程低延迟交换数据。
- 寄存器文件（Register File）：每个SM拥有庞大的寄存器堆（如256KB），用于存放线程的局部变量。寄存器是GPU中最快的存储，但总量有限，直接限制可并发的线程数。
- 线程束调度器（Warp Scheduler）：每个SM包含4~6个调度器，每个调度器负责管理一组线程束（Warp，通常为32个线程），以“锁步”方式分发指令。
内存层次是GPU架构的另一核心。从寄存器到全局显存，延迟和带宽差异巨大，理解这一层次结构对优化访存模式至关重要。
| 存储层级 | 位置 | 典型容量 | 延迟（周期） | 作用域 | 备注 |
|---|---|---|---|---|---|
| 寄存器文件 | SM内部 | 256KB/SM | ~1 | 单个线程 | 最快，但需避免溢出 |
| 共享内存 | SM内部 | 48~228KB/SM | ~20 | 线程块内 | 可编程，需显式管理 |
| L1缓存 | SM内部 | 与共享内存共享 | ~30 | 线程块内 | 自动缓存全局内存访问 |
| L2缓存 | 芯片级 | 4~40MB | ~200 | 全局 | 所有SM共享 |
| 全局显存（HBM/GDDR） | 片外 | 16~80GB | ~400-800 | 所有线程 | 容量大，带宽高，延迟极高 |
下图展示了SM内部结构及线程调度关系，帮助建立直观的硬件认知：
并行计算模型：从线程到线程束
CUDA采用单指令多线程（SIMT）执行模型。程序员编写的核函数（Kernel）由网格（Grid）启动，网格由若干线程块（Thread Block）组成，每个线程块包含若干线程（Thread）。硬件执行时，线程块被分配到SM上，SM再将线程块内的线程划分为线程束（Warp）——每个Warp固定包含32个线程，这是硬件调度的最小单位。
Warp调度器每个周期选择一个就绪的Warp，将其下一条指令广播给该Warp内所有线程执行。这意味着：
- 同一Warp内的32个线程必须执行相同指令。若出现分支发散（如if-else），则不同路径会被串行执行，导致部分线程闲置，性能下降。
- 当Warp内某个线程访问全局内存时，硬件会尝试将32个线程的访存请求合并（Coalescing）为尽可能少的内存事务。若地址连续，则一次事务即可完成；若地址分散，则需多次事务，浪费带宽。
- SM通过线程块级并发隐藏访存延迟：当一个Warp因等待内存数据而阻塞时，调度器立即切换到另一个就绪Warp执行计算，无需等待。
实践要点：核函数配置（Grid/Block维度）直接影响性能。通常建议Block大小设为32的倍数（如128、256），以保证Warp完整填充。Block过大可能导致寄存器溢出，过小则无法充分隐藏延迟。需结合具体硬件资源（寄存器数、共享内存上限）综合权衡。
案例：矩阵乘法的并行化映射
以经典的C = A × B（矩阵尺寸1024×1024）为例，展示并行计算模型如何落地。采用分块（Tiling）策略：
// 核函数：每个线程计算C中一个元素
__global__ void matmul_kernel(float *A, float *B, float *C, int N) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int k = 0; k < N; ++k) {
sum += A[row * N + k] * B[k * N + col];
}
C[row * N + col] = sum;
}
// 启动配置：Block=16×16=256线程，Grid=64×64=4096个Block
dim3 block(16, 16);
dim3 grid(N/16, N/16);
matmul_kernel<<<grid, block>>>(A, B, C, N);该朴素实现存在严重的访存效率问题：每个线程直接读取全局内存中的A和B，且相邻线程访问的B列地址不连续，导致访存无法合并。优化方案是使用共享内存分块：
- 将A和B分别划分为16×16的小块（Tile），与Block尺寸对应。
- 每个Block协作将所需的A块和B块从全局内存加载到共享内存。
- 线程从共享内存读取数据计算，大幅减少全局内存访问次数。
- 循环遍历所有Tile，累加部分和。
优化后，每个线程对全局内存的访问从2×1024次降至2×(1024/16)=128次，访存效率提升16倍。此案例揭示了GPU优化的核心矛盾：算力远高于显存带宽，因此减少全局内存访问往往比减少计算量更关键。
误区与复盘：常见性能陷阱
误区一：盲目增加线程数就能提升性能。线程过多会导致每个线程占用的寄存器减少，可能触发寄存器溢出（Spill）到局部内存（实际在显存中），反而大幅降低性能。应通过Occupancy API（如cudaOccupancyMaxActiveBlocksPerMultiprocessor）计算理论占用率，找到平衡点。
误区二：忽视Warp发散。在核函数中写if (threadIdx.x % 2 == 0)这类分支，会导致Warp内一半线程执行路径A、另一半执行路径B，最终串行执行两条路径，吞吐量减半。应尽量让分支条件基于Warp粒度（如warp_id % 2），或重构算法避免发散。
误区三：共享内存使用不当。共享内存虽快，但容量有限且需要__syncthreads()同步。若每个Block使用超过SM可用共享内存，核函数将无法启动。此外，共享内存Bank冲突（多个线程同时访问同一Bank的不同地址）会导致访问串行化，需通过填充（Padding）或调整访问模式规避。
复盘一个真实案例：某图像处理核函数在V100上实测占用率仅25%，经ncu分析发现寄存器使用量高达128个/线程，导致每个SM仅能驻留4个Block（2048线程上限的1/4）。通过将部分中间变量改为__shared__数组，寄存器降至64个，占用率提升至50%，性能提升1.8倍。这印证了资源占用与并行度的直接权衡——寄存器、共享内存、线程数三者构成“不可能三角”，需根据核函数特性动态调整。