2、GPU硬件架构基础：SM、CUDA核心与内存层次结构

要真正掌握CUDA编程，不能只停留在“写一个kernel函数”的层面。你必须理解GPU硬件是如何组织计算单元、如何调度线程、以及数据在各级存储之间如何流动。本章将从NVIDIA GPU的物理架构出发，拆解流式多处理器（SM）、CUDA核心、内存层次结构以及GPU的并行计算模型，帮助你建立从硬件到编程模型的映射关系。只有理解了硬件，你才能解释为什么某些代码快、某些代码慢，也才能写出真正发挥GPU性能的程序。

2.1 流式多处理器（SM）：GPU的基本计算单元

SM（Streaming Multiprocessor）是GPU物理架构中的核心计算单元。一个GPU由多个SM组成，例如NVIDIA A100拥有108个SM，而消费级RTX 3060则有28个SM。每个SM内部包含若干CUDA核心、共享内存、寄存器文件、缓存以及调度器。SM的设计哲学是“小而密”——单个SM的计算能力有限，但通过大量SM并行工作，GPU实现了极高的吞吐量。

在编程模型层面，CUDA将线程组织为线程块（Thread Block），而硬件调度器会将每个线程块分配给一个SM执行。一个SM可以同时驻留多个线程块，具体数量取决于线程块占用的资源（共享内存、寄存器等）。SM内部的线程调度以Warp为单位——每个Warp包含32个线程，这是硬件执行的基本粒度。GPU采用SIMT（Single Instruction, Multiple Thread）模型，即一个Warp中的所有线程在同一时刻执行同一条指令，但可以处理不同的数据。

2.2 CUDA核心：并非传统CPU核心

CUDA核心（CUDA Core）是SM内部最基本的计算单元，负责执行浮点运算（FP32）、整数运算（INT32）等算术操作。以A100为例，每个SM包含64个FP32 CUDA核心和64个INT32核心。需要注意的是，CUDA核心远不如CPU核心“智能”——它没有分支预测、乱序执行等复杂逻辑，而是依赖大量核心并行来弥补单核能力的不足。

在编程时，你不需要直接控制CUDA核心，而是通过编写kernel函数，让GPU自动将线程映射到这些核心上。每个线程执行kernel函数中的代码，而硬件调度器负责将Warp分发到CUDA核心上执行。一个Warp中的32个线程会被分配到32个CUDA核心上并行执行（如果SM有足够的核心）。

2.3 内存层次结构：从寄存器到全局内存

GPU的内存层次结构是性能调优的核心。从快到慢、从容量小到大，依次为：寄存器、共享内存、L1/L2缓存、全局内存、常量内存和纹理内存。理解每一级内存的延迟、带宽和用途，是编写高效CUDA程序的前提。

2.3.1 寄存器（Register）

寄存器是速度最快、延迟最低的存储，位于每个SM内部，以寄存器文件的形式存在。每个线程的局部变量（如循环索引、临时计算值）默认存储在寄存器中。A100每个SM有256 KB寄存器文件，即65536个32位寄存器。每个线程最多可以使用255个寄存器，但SM的总寄存器数量有限，因此线程使用的寄存器越多，SM能同时驻留的线程就越少。

2.3.2 共享内存（Shared Memory）

共享内存是SM内部的一块可编程高速缓存，供同一个线程块内的所有线程访问。它的延迟约为20~30ns，远低于全局内存的400~600ns。共享内存常用于线程块内的数据复用、协作计算和避免重复访问全局内存。A100每个SM有164 KB共享内存，可以通过__shared__关键字在kernel中声明。

2.3.3 全局内存（Global Memory）

全局内存是GPU中容量最大但延迟最高的存储，通常由HBM2e或GDDR6显存构成。所有线程都可以访问全局内存，但必须通过L2缓存。全局内存的带宽虽然很高（A100约2 TB/s），但延迟极大，因此应尽量减少全局内存的访问次数，或者通过合并访问（Coalesced Access）提高带宽利用率。

2.3.4 常量内存与纹理内存

常量内存（Constant Memory）是只读的，容量为64 KB，适合存储所有线程共享的常量数据。当所有线程访问同一地址时，常量内存的广播机制非常高效。纹理内存（Texture Memory）也是只读的，专为图像处理设计，支持硬件插值和边界处理，适合二维空间数据的访问模式。

2.4 GPU并行计算模型：从线程到Warp再到Grid

CUDA的并行计算模型是分层的：Grid → Block → Thread → Warp。当你启动一个kernel时，实际上创建了一个Grid，Grid由多个Block组成，每个Block包含多个Thread。硬件将Block调度到SM上，SM再将Thread划分为Warp（每32个线程一组）进行调度。

1. Grid：一次kernel启动的所有线程集合，由gridDim.x/y/z定义。
2. Block（线程块）：Grid的子集，由blockDim.x/y/z定义。Block内的线程可以通过共享内存通信，并通过__syncthreads()同步。
3. Thread（线程）：执行kernel函数的最小单位，通过threadIdx.x/y/z索引。
4. Warp：硬件调度的基本单位，32个线程组成一个Warp。Warp中的线程执行相同的指令，但处理不同的数据。

这种分层模型的关键在于：Block内的线程可以协作，而Block之间不能直接通信（只能通过全局内存间接通信）。因此，在设计kernel时，应将需要协作的线程放在同一个Block内，而将相互独立的任务分配到不同的Block上。

2.5 案例拆解：矩阵乘法的硬件映射

以经典的矩阵乘法C = A × B为例，假设矩阵大小为1024×1024。我们设计一个kernel，每个线程计算C中的一个元素。线程块大小为16×16=256个线程，Grid大小为64×64=4096个Block。

__global__ void matMul(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线程）被调度到一个SM上。SM将256个线程划分为8个Warp（每个Warp 32线程）。每个Warp中的线程访问A和B的相邻元素，这有利于合并内存访问。然而，这个朴素版本存在严重的性能问题：每个线程需要读取A的一行和B的一列，导致大量重复的全局内存访问。

优化方案：使用共享内存分块（Tiling）。将矩阵划分为16×16的块，每个Block先将A和B的对应块加载到共享内存，然后进行块内计算。这样，每个全局内存元素只被读取一次，大幅减少了全局内存访问次数。

2.6 常见误区与调优要点

• 误区1：线程越多越好。实际上，过多的线程会导致资源不足（寄存器、共享内存耗尽），反而降低SM驻留的Warp数量，影响延迟隐藏能力。
• 误区2：忽略Warp发散。如果Warp中的线程执行不同的分支（if-else），GPU会串行执行每个分支，导致性能下降。应尽量让Warp内的线程走相同的执行路径。
• 误区3：全局内存随机访问。全局内存的带宽利用率依赖于合并访问——即同一Warp的32个线程访问连续的内存地址。随机访问会导致多次内存事务，严重降低带宽。
• 误区4：共享内存使用不当。共享内存虽然快，但容量有限（164 KB/SM），且存在bank conflict问题。当多个线程同时访问同一bank的不同地址时，访问会被串行化。

2.7 总结复盘