GPU架构基础：SM、线程层次、内存层次与SIMT执行模型

CUDA编程模型之所以能够将GPU的海量并行计算能力释放出来，其底层依赖的是一套精心设计的硬件架构与执行逻辑。理解GPU的物理架构（SM）、逻辑组织（线程层次）、数据驻留（内存层次）以及指令执行方式（SIMT），是编写高性能cuBLAS替代实现或自定义CUDA内核的基石。本章将深入剖析这四大核心要素，帮助你在GPU上建立正确的“硬件心智模型”，从而在后续的矩阵运算优化中做出合理的决策。

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

核心原理

SM（Streaming Multiprocessor）是GPU物理硬件的基本计算单元，类似于CPU中的多核处理器核心，但设计哲学截然不同。一个现代GPU（如NVIDIA A100）包含108个SM，每个SM内部又包含多个CUDA核心（或称SP，Streaming Processor）、共享内存、寄存器堆、特殊功能单元（SFU）、加载/存储单元（LD/ST）以及调度器。SM采用单指令多线程（SIMT）架构，即同一时刻，SM中的所有执行单元会执行同一条指令，但作用于不同的数据。

关键点在于，SM是资源分配和调度的物理边界。当一个CUDA内核启动时，线程块（Block）会被分配到SM上执行。一个SM可以同时驻留多个线程块，具体数量取决于线程块所需的资源（共享内存、寄存器）与SM可用资源的比值。SM内部的线程调度器（Warp Scheduler）负责将线程块中的线程进一步划分为更小的调度单位——Warp。

判断方法与操作流程

在编写或优化CUDA程序时，你需要通过以下方式感知SM的存在与限制：

1. 查询设备属性：使用 cudaGetDeviceProperties() 获取 multiProcessorCount（SM数量）、maxThreadsPerMultiProcessor（每个SM最大线程数）、sharedMemPerMultiprocessor（每个SM共享内存大小）等关键参数。
2. 计算占用率（Occupancy）：占用率 = (活动Warp数 / 最大Warp数) × 100%。高占用率有助于隐藏内存延迟，但并非越高越好，过高的占用率可能导致寄存器溢出（Spill）到本地内存，反而降低性能。
3. 资源分配决策：在启动内核时，通过执行配置（Execution Configuration）指定线程块大小和网格大小。你需要根据SM的资源容量，合理选择每个线程块的大小（通常为128或256的倍数），以最大化SM的资源利用率。

案例拆解：矩阵乘法中的SM资源分配

假设我们要在A100上执行一个 1024×1024 的矩阵乘法。A100每个SM有64K个32位寄存器，最大支持2048个线程驻留。如果我们设置线程块大小为256，那么每个SM最多可以驻留 2048/256 = 8 个线程块。如果每个线程需要32个寄存器，那么每个线程块需要 256×32 = 8192 个寄存器，8个线程块共需 65536 个寄存器，恰好占满整个SM的寄存器堆。此时占用率达到100%。

但如果我们将线程块大小改为512，每个SM最多驻留4个线程块，总线程数仍为2048，占用率不变。然而，如果每个线程的寄存器使用量增加到64个，那么每个线程块需要 512×64 = 32768 个寄存器，4个线程块共需 131072 个寄存器，远超SM的64K寄存器容量。此时，编译器会将多余的寄存器变量溢出到本地内存（Local Memory），导致额外的内存访问开销，性能急剧下降。因此，在优化时，我们需要通过 __launch_bounds__ 或 maxrregcount 控制寄存器使用量，以平衡占用率和单线程性能。

2.2 线程层次：Grid、Block与Warp

核心原理

CUDA的线程组织采用三层结构：Grid（网格）→ Block（线程块）→ Thread（线程）。这种层次结构是逻辑上的，它允许程序员以灵活的方式组织并行任务，同时映射到物理硬件上。一个Grid包含多个Block，一个Block包含多个Thread。Block内的线程可以通过共享内存和同步屏障（__syncthreads()）进行协作，而Block之间则相互独立，无法直接通信或同步。

在硬件执行层面，Block会被调度到SM上，而SM内部将Block进一步划分为Warp。Warp是GPU执行指令的最小调度单位，通常包含32个线程。这意味着，无论Block中有多少个线程，硬件总是以32个线程为一组来发射和执行指令。如果Block大小不是32的倍数，最后一个Warp中会有部分线程处于空闲状态（Inactive），造成计算资源浪费。

判断方法与操作流程

选择合适的线程层次结构是CUDA编程的核心技能之一，通常遵循以下流程：

• 确定Block大小：一般建议设置为32的倍数（如128、256、512）。较小的Block（如64）有利于提高占用率，但可能增加调度开销；较大的Block（如1024）适合计算密集型任务，但可能降低SM上驻留的Block数量。
• 确定Grid大小：Grid大小 = 总线程数 / Block大小。总线程数应覆盖所有需要处理的数据元素。例如，处理1000个元素，Block大小为256，则需要4个Block（最后一个Block有24个线程空闲）。
• 使用内置变量：在核函数中，通过 blockIdx.x、blockDim.x、threadIdx.x 计算全局线程ID：int tid = blockIdx.x * blockDim.x + threadIdx.x;。

案例拆解：向量加法的线程映射

考虑一个长度为 N=10000 的向量加法。我们设置Block大小为256，则Grid大小为 ceil(10000/256) = 40。在核函数中，每个线程处理一个元素：

__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

这里，if (i < N) 是边界检查，防止越界访问。在40个Block中，前39个Block的256个线程全部活跃，最后一个Block只有 10000 - 39×256 = 16 个线程活跃，其余240个线程空闲。这种“尾部效应”在数据规模不是Block大小整数倍时不可避免，但可以通过调整Block大小（如使用128）来减少空闲线程比例。

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

核心原理

GPU的内存层次结构是决定性能的关键因素之一。从快到慢、从近到远依次为：寄存器（Register）→ 共享内存（Shared Memory）→ L1/L2缓存 → 全局内存（Global Memory）→ 本地内存（Local Memory）。此外，还有只读的常量内存（Constant Memory）和纹理内存（Texture Memory）。

寄存器和共享内存是片上资源，访问延迟极低（约1-30个时钟周期），但容量有限。全局内存是显存（DRAM），容量大但延迟高（约400-800个时钟周期）。因此，性能优化的核心原则是：最大化数据在寄存器/共享内存中的复用，最小化对全局内存的访问次数。

判断方法与操作流程

在优化内存访问时，可以遵循以下决策流程：

1. 优先使用寄存器：将频繁访问的循环变量、中间计算结果存储在局部变量中，编译器会自动将其分配到寄存器。
2. 利用共享内存做数据复用：当Block内的多个线程需要访问同一份数据（如矩阵乘法中的子矩阵块）时，先将数据从全局内存加载到共享内存，再进行计算。
3. 保证全局内存的合并访问（Coalesced Access）：同一Warp中的32个线程，如果访问的全局内存地址是连续的，硬件会将这32次访问合并为一次或少数几次内存事务，极大提高带宽利用率。
4. 避免本地内存溢出：如果每个线程使用的寄存器超过硬件限制，编译器会将多余的变量放入本地内存（实际是全局内存），导致性能骤降。可以通过 -Xptxas -v 编译选项查看寄存器使用情况。

案例拆解：矩阵转置中的内存访问优化

矩阵转置是一个经典的内存访问优化案例。朴素实现中，每个线程读取一行写入一列，导致写入操作无法合并访问。优化方案是使用共享内存作为中转：

__global__ void transpose(float *in, float *out, int n) {
    __shared__ float tile[TILE][TILE];
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;
    if (x < n && y < n) {
        tile[threadIdx.y][threadIdx.x] = in[y * n + x];
    }
    __syncthreads();
    int x2 = blockIdx.y * TILE + threadIdx.x;
    int y2 = blockIdx.x * TILE + threadIdx.y;
    if (x2 < n && y2 < n) {
        out[y2 * n + x2] = tile[threadIdx.x][threadIdx.y];
    }
}

在这个实现中，tile 是共享内存中的二维数组。线程首先协作将全局内存中的一块数据加载到共享内存（合并读取），同步后，再从共享内存中读取并写入全局内存（合并写入）。通过共享内存的“转置”操作，避免了全局内存的非合并访问，性能可提升数倍。

2.4 SIMT执行模型：Warp的并行与分歧

核心原理

SIMT（Single Instruction, Multiple Thread）是GPU的执行模型。与CPU的SIMD（单指令多数据）不同，SIMT允许每个线程拥有独立的寄存器、独立的指令地址和独立的数据路径。然而，硬件在执行时，是以Warp（32个线程）为单位，由Warp调度器向所有活跃线程发射同一条指令。这意味着，同一Warp中的线程必须执行相同的指令，如果出现分支分歧（Branch Divergence），即不同线程走不同的代码路径，硬件会串行执行每个分支路径，并屏蔽不活跃的线程，导致性能下降。

判断方法与操作流程

避免分支分歧是SIMT编程的重要技巧，具体方法包括：

• 将分支条件转换为算术运算：例如，使用 a = (condition) ? x : y; 代替 if (condition) a = x; else a = y;。
• 重新组织数据：将需要执行相同操作的数据分组，使Warp内的线程尽可能执行相同的指令。
• 使用谓词（Predication）：对于简单的条件赋值，编译器会自动使用谓词指令，避免分支。
• 注意循环边界：当循环次数不是Warp大小的整数倍时，最后一个Warp中的部分线程会提前退出循环，造成资源浪费。

案例拆解：粒子碰撞检测中的分支优化

假设一个粒子模拟程序，需要检测粒子是否与边界碰撞。朴素实现中，每个线程处理一个粒子，并根据粒子的位置执行不同的碰撞响应代码：

if (x < 0) { x = -x; vx = -vx; }
else if (x > width) { x = 2*width - x; vx = -vx; }
// 其他边界...

如果粒子随机分布，那么在一个Warp中，可能有的线程执行 x < 0 分支，有的执行 x > width 分支，有的不执行任何分支。这会导致Warp调度器串行执行三个分支路径，总执行时间约为单个分支的三倍。

优化方案是使用无分支的数学表达式：

// 使用 clamp 和符号翻转
bool left = (x < 0);
bool right = (x > width);
x = left ? -x : (right ? 2*width - x : x);
vx = (left || right) ? -vx : vx;

虽然这里仍然使用了三元运算符，但现代编译器会将其编译为谓词指令（如 FSEL），而不是真正的分支跳转。这样，所有线程都执行相同的指令序列，只是根据条件选择不同的操作数，从而避免了分支分歧。

2.5 总结与性能调优要点

本章从硬件架构到执行模型，系统梳理了GPU并行计算的基础。核心要点归纳如下：

在后续的cuBLAS库分析与自定义内核开发中，你将反复运用这些原则。例如，cuBLAS的矩阵乘法内核（如 cublasSgemm）之所以性能卓越，正是因为它精心设计了线程块大小（通常为128或256）、共享内存分块（如16×16或32×32的