GPU内存基础：全局内存、共享内存、寄存器、局部内存的特性与访问成本

在深入优化GPU内核之前，必须先建立一张清晰的内存层级地图。GPU并非拥有单一均匀的存储池，而是由多种物理与逻辑上分离的内存空间构成，每种内存的容量、速度、作用域和编程模型都截然不同。理解这些差异，是判断一个内核为何慢、以及如何改快的起点。

核心原理：四种内存的定位与硬件映射

CUDA（以及HIP等同类模型）将可编程内存抽象为四个主要层级：寄存器、共享内存、全局内存和局部内存。它们并非并列关系，而是分布在从芯片内部到板卡显存的不同物理位置。

• 寄存器（Register）：位于SM（流式多处理器）内部，是速度最快的存储单元，零额外延迟。每个线程独占其寄存器，生命周期与线程一致。寄存器由编译器静态分配，若一个线程需要的寄存器数量超过硬件上限，就会发生“寄存器溢出”。
• 共享内存（Shared Memory）：同样位于SM内部，但由线程块（Block）内所有线程共享。它是一块可编程管理的片上SRAM，延迟远低于全局内存，但容量有限（现代GPU通常为每SM 64KB~228KB，可配置分区）。共享内存需要程序员显式声明和同步访问。
• 全局内存（Global Memory）：位于显存（DRAM）中，是所有线程（包括主机端）都能访问的庞大存储池。容量可达数GB至数十GB，但访问延迟高达数百个时钟周期。全局内存通过L2缓存和L1缓存（部分架构）进行加速，但本质上属于片外存储。
• 局部内存（Local Memory）：名字具有迷惑性——它并不在芯片上，而是全局内存中为每个线程私有的一部分逻辑空间。当寄存器溢出、或访问动态索引的局部数组时，编译器会将数据“降级”到局部内存。其访问成本与全局内存相当，且会失去缓存行合并的优势。

访问成本与带宽特性对比

下表汇总了四种内存的关键特性，这是后续所有优化决策的量化基础。请注意，具体数值因GPU架构（如Ampere、Hopper、Ada Lovelace）而异，但量级关系保持不变。

内存类型	物理位置	访问延迟（周期）	容量/线程	作用域	缓存策略	典型用途
寄存器	SM内	~1（无延迟）	最多255个（32位）	单线程	无	循环变量、中间计算结果
共享内存	SM内	~20-30	最多48KB/Block（可申请更多）	线程块	无（手动管理）	数据重用、块内通信、广播
全局内存	板卡DRAM	~400-800	数GB~数十GB	所有线程+主机	L1/L2	输入数据、输出结果、持久化数组
局部内存	板卡DRAM	~400-800	受限于显存总量	单线程（逻辑私有）	L1/L2（但无合并保证）	寄存器溢出、动态索引的局部数组

方法流程：如何诊断内存使用特征

在动手优化前，必须通过工具量化内核的内存行为。推荐使用NVIDIA Nsight Compute（或AMD的rocprof）进行以下三步诊断：

1. 查看寄存器与局部内存统计：在“Memory Workload Analysis”页面中，检查“Register Usage”和“Local Memory”指标。若局部内存字节数不为0，需检查溢出原因。
2. 分析全局内存访问模式：关注“Memory Throughput”和“Sectors/Request”。如果每个请求的扇区数远大于理论最小值，说明存在非合并访问。
3. 检查共享内存Bank冲突：在“Shared Memory”分析中，查看“Bank Conflicts”计数。高冲突率会显著降低共享内存的有效带宽。

以下是一个典型的诊断命令示例（Nsight Compute）：

# 分析内核的内存工作负载
ncu --section MemoryWorkloadAnalysis --section SpeedOfLight ./my_kernel

# 输出关键指标示例：
# Register Usage: 32 registers/thread
# Local Memory: 8 bytes/thread (溢出!)
# Global Load Efficiency: 50% (存在非合并访问)
# Shared Memory Bank Conflicts: 1024 (高冲突)

案例/场景：矩阵乘法中的内存角色分配

以经典的SGEMM（单精度矩阵乘法）为例，展示四种内存如何协同工作。假设计算 C = A × B，其中矩阵尺寸为 N×N。

朴素实现（仅用全局内存）：每个线程直接读取A和B的全局内存元素。由于A和B的元素会被多个线程重复使用，每次读取都触发全局内存访问，导致极高的内存流量。例如，计算一个 32×32 的输出Tile，需要读取 32×32×2 个浮点数，但实际有效数据仅为 32×32 个，重复读取率高达32倍。

优化实现（分块+共享内存）：

1. 全局内存→共享内存：每个线程块协作，将A和B的一个 32×32 Tile从全局内存加载到共享内存。此过程需保证合并访问——相邻线程读取相邻地址。
2. 共享内存→寄存器：每个线程从共享内存中读取所需的A和B元素到寄存器中。由于共享内存延迟低，且可通过广播避免Bank冲突，此步骤开销极小。
3. 寄存器计算：在寄存器中执行乘加运算，结果累积在寄存器中，避免频繁访问共享内存。
4. 写回全局内存：计算完成后，将寄存器中的结果Tile写回全局内存。

该优化将全局内存访问次数从 O(N³) 降低到 O(N²)，性能提升可达10倍以上。核心思想是：用共享内存的显式管理，换取全局内存流量的指数级下降。

误区与复盘：局部内存的“隐形陷阱”

以下是一个真实案例，展示了局部内存如何悄然破坏性能。某内核代码中，每个线程声明了一个长度为16的局部数组，并希望通过索引变量动态访问：

__global__ void process(float* data) {
    int idx = threadIdx.x % 4;  // 动态索引
    float temp[16];             // 局部数组
    for (int i = 0; i < 16; ++i) {
        temp[i] = data[i * blockDim.x + threadIdx.x];
    }
    // 动态访问temp
    float result = temp[idx * 4 + 1];
    // ... 后续计算
}

问题分析：由于 temp 的索引 idx*4+1 在编译期无法确定，编译器无法将 temp 完全映射到寄存器，只能将其放置在局部内存中。每个线程访问 temp 时，都会产生一次全局内存级别的读写。更糟的是，这些访问是分散的（每个线程的 temp 地址不同），无法合并，导致L1缓存命中率极低。

性能数据：使用Nsight Compute分析发现，该内核的局部内存流量高达 2.5GB/s，而全局内存有效吞吐仅为 0.8GB/s。局部内存访问占据了约70%的内存指令周期。

修复方案：将 temp 数组改为固定索引访问，或使用 #pragma unroll 强制循环展开，使编译器能够将数组元素分配到独立寄存器中。修改后，局部内存使用降为0字节，内核性能提升3.2倍。

上图直观展示了四种内存的物理位置与访问路径。注意局部内存的虚线框——它虽然逻辑上属于线程私有，但物理上位于显存区域，与全局内存共享同一DRAM带宽。这也是为什么局部内存访问过多会同时拖累全局内存性能的原因。

在实际优化中，建议遵循以下优先级原则：寄存器 > 共享内存 > 全局内存（合并访问） > 局部内存（尽量避免）。当寄存器不足时，优先考虑通过算法重构减少寄存器压力，而非依赖局部内存；当共享内存不足时，考虑分块（tiling）策略或调整块大小，而非直接回退到全局内存。