📚 课程目录第4章 GPU内存模型

GPU内存模型：全局内存、共享内存、寄存器与缓存层次

GPU 的高吞吐量不仅来自数千个计算核心，更依赖一套层次分明、带宽差异巨大的内存体系。理解这一模型，是编写高性能 CUDA/HIP 内核的前提。本章从硬件视角拆解各级内存的物理位置、容量、延迟与用途，并给出可落地的优化策略。

4.1 内存层次全景：从寄存器到全局内存

下图展示了 GPU 内存体系的完整结构。从最靠近计算单元的寄存器，到最远的全局内存（显存），每一级都有明确的角色与性能边界。

上图中，寄存器是每个线程私有的存储，速度最快但容量极小（每线程通常最多 255 个 32 位寄存器）。共享内存（Shared Memory）位于 SM 内部，同一线程块内的线程可以互相访问，延迟远低于全局内存。L1/L2 缓存对程序员部分透明，但理解其行为对优化至关重要。全局内存则是所有线程共享的大容量存储，也是 CPU 与 GPU 数据交换的主通道。

4.2 核心原理：为什么内存模型决定性能上限

GPU 性能瓶颈往往不在计算，而在数据搬运。一个典型的算术密集型内核，若每个浮点运算需要从全局内存读取操作数，其实际吞吐量可能仅为峰值算力的 5%-10%。因此，优化内存访问模式比优化指令序列更重要。

三个核心原理需要牢记：

合并访问（Coalesced Access）：同一 warp（32 线程）访问连续地址时，硬件将其合并为少数几次内存事务。若访问跳跃或分散，事务数增加，有效带宽骤降。
局部性（Locality）：时间局部性（同一数据多次使用）和空间局部性（相邻数据同时使用）决定了缓存命中率。利用共享内存或寄存器缓存复用数据，可显著减少全局内存访问。
占用率（Occupancy）：每个 SM 能同时驻留的线程数受寄存器、共享内存等资源限制。资源占用过高会降低占用率，导致无法用线程切换隐藏内存延迟。

经验法则： 全局内存每次访问的代价约为几百个周期。若一个线程块需要反复读取同一份数据，应先将数据加载到共享内存，再通过 __syncthreads() 同步后使用。这能将数百次全局访问降为一次。

4.3 优化方法流程：从分析到落地

内存优化不是盲目套用技巧，而是系统性的分析过程。以下流程适用于大多数内核：

Profiler 定位瓶颈：使用 ncu（NVIDIA Nsight Compute）或 rocprof（AMD）分析内核，查看 Memory Throughput 和 Compute Throughput。若前者接近 100% 而后者较低，则属于内存受限。
检查访问模式：审查全局内存访问是否合并。例如，以下代码实现了合并访问：

// 合并访问：线程 i 访问 data[i]
__global__ void copy_kernel(float *in, float *out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i];  // 相邻线程访问相邻地址
}

// 非合并访问：线程 i 访问 data[i * 32]（跨步访问）
__global__ void stride_kernel(float *in, float *out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i * 32];  // 每次事务仅利用 1/32 数据
}

数据复用分析：若每个数据被同一线程块内多个线程使用，考虑用共享内存做块内缓存。典型场景是矩阵乘法中的分块（tiling）。
调整资源分配：通过 cudaFuncSetAttribute() 或启动配置调整共享内存/寄存器分配，平衡占用率与单线程性能。
迭代验证：每次修改后重新 profiling，对比带宽与耗时变化。

4.4 案例：矩阵乘法分块优化

以 1024×1024 的矩阵乘法 C = A × B 为例。朴素实现中，每个线程直接读取全局内存中的 A 和 B，每个输出元素需要 2N 次全局访问。分块后，每个线程块计算一个 16×16 的输出子块，并将对应的 A 子块（16×1024）和 B 子块（1024×16）加载到共享内存。

#define TILE_SIZE 16

__global__ void matmul_tiled(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    for (int tile = 0; tile < N / TILE_SIZE; ++tile) {
        // 协作加载到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * N + col];
        __syncthreads();

        for (int k = 0; k < TILE_SIZE; ++k)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();  // 防止下一轮覆盖前数据
    }
    C[row * N + col] = sum;
}

该优化将全局内存访问量从 O(N³) 降至 O(N³ / TILE_SIZE)，实际加速比可达 5-10 倍（取决于矩阵规模与硬件）。关键点在于：每个线程块只加载自己需要的子块，且通过 __syncthreads() 保证数据就绪后再计算。

4.5 误区与复盘：常见性能陷阱

误区	表现	正确做法
忽略合并访问	带宽利用率低于 20%	确保 threadIdx.x 映射到连续内存维度
共享内存使用过度	占用率下降，延迟无法隐藏	平衡共享内存大小与 block 数量
未使用只读缓存	纹理/常量数据反复访问全局	使用 `__ldg()` 或 const __restrict__ 提示
同步过多	线程等待时间 > 计算时间	减少同步点，或使用异步拷贝（cp.async）

典型复盘：某图像处理内核将每个像素的 3×3 邻域数据重复从全局内存读取，导致 9 倍冗余访问。改为共享内存后，每个线程块只加载一次图像块（含 halo 边界），性能提升 4.2 倍。教训：先分析数据复用模式，再决定缓存策略，不要盲目增加共享内存大小。

4.6 关键参数速查

不同架构的 GPU 内存参数略有差异，以下以 NVIDIA H100（Hopper）和 A100（Ampere）为例：

寄存器文件：每 SM 256KB，每线程最多 255 个寄存器。
共享内存：每 SM 最大 228KB（H100），可配置为 L1 缓存拆分。
L2 缓存：H100 为 50MB，A100 为 40MB。
全局内存带宽：H100 SXM 约 3.35 TB/s，A100 约 2.0 TB/s。
缓存行大小：全局内存访问以 32 字节扇区为单位，合并访问应尽量覆盖完整扇区。

实际开发中，建议先用 cudaDeviceProp 查询当前设备的 sharedMemPerBlock、regsPerMultiprocessor 等参数，再据此设计内核配置。