GPU内存层次结构详解

大家好，我是你们的CUDA讲师。今天我们来聊聊GPU的内存体系——说实话，这是整个CUDA编程里最绕不开的核心话题。我见过太多同学，kernel写得挺漂亮，一跑性能就是上不去，最后发现是内存访问模式出了问题。

GPU的内存不像CPU那样简单。它是一套层次化的结构，每一层都有不同的延迟、容量和用途。你想想看，不理解这个层次结构，就像开着一辆超跑却不知道换挡——再强的算力也发挥不出来。

全局内存（Global Memory）

全局内存是GPU里容量最大的存储空间，通常有数GB到数十GB。所有线程都能访问它，但代价是——延迟非常高。

我个人习惯把全局内存比作「仓库」。东西多，但取一次货要走很远。我在项目中遇到过这样一个案例：一个矩阵乘法kernel，最初直接读写全局内存，性能只有理论峰值的15%。后来改成共享内存分块，直接飙到70%以上。嗯，差距就是这么明显。

访问全局内存时，有个非常重要的原则——合并访问（Coalesced Access）。说白了就是：相邻的线程要访问相邻的地址。为什么？因为GPU会把多个线程的访存请求合并成一次大的内存事务。

// 好的访问模式：合并访问
__global__ void coalesced_kernel(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[idx];  // 相邻线程访问相邻地址
}

// 坏的访问模式：非合并访问
__global__ void strided_kernel(float* data) {
    int idx = threadIdx.x * 32 + blockIdx.x * blockDim.x;
    float val = data[idx];  // 相邻线程访问间隔32的地址
}

共享内存（Shared Memory）

共享内存是GPU里最「香」的存储资源之一。它位于芯片内部，延迟极低——大约只有全局内存的1/10到1/20。

属性	值
容量	每个SM通常48KB~164KB（可配置）
延迟	约20~40个时钟周期
作用域	同一个线程块内的所有线程
生命周期	线程块执行期间

共享内存的本质是什么？是程序员可控的缓存。你可以手动把数据从全局内存搬到共享内存，然后让线程块内的所有线程快速访问。这特别适合数据重用场景——比如卷积、矩阵乘法。

__global__ void shared_mem_example(float* input, float* output, int N) {
    __shared__ float s_data[256];  // 声明共享内存
    
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;
    
    // 从全局内存加载到共享内存
    if (gid < N) {
        s_data[tid] = input[gid];
    }
    __syncthreads();  // 同步！确保所有线程都加载完毕
    
    // 在共享内存上做计算
    // ... 这里可以安全地访问 s_data
    
    // 写回全局内存
    if (gid < N) {
        output[gid] = s_data[tid] * 2.0f;
    }
}

寄存器（Register）

寄存器是GPU里最快的存储，没有之一。每个线程都有自己的私有寄存器，延迟只有1个时钟周期。

但寄存器有个大问题——数量有限。每个SM的寄存器总数是固定的（比如65536个），平均分给所有活跃线程。如果你每个线程用太多寄存器，就会降低SM上同时运行的线程数（occupancy）。

怎么查看寄存器使用量？用--ptxas-options=-v编译选项，或者用NVIDIA Nsight工具。我个人习惯在开发阶段经常检查这个指标。

本地内存（Local Memory）

本地内存这个名字容易让人误解。它其实不是「本地」的——它存储在全局内存里，只是每个线程私有的。当寄存器不够用时，编译器会自动把变量「溢出」到本地内存。

怎么避免？减少每个线程的变量数量，或者用共享内存替代部分局部变量。

常量内存（Constant Memory）

常量内存是GPU里一个特殊的存在。它容量很小（通常64KB），但有两个优势：

• 有专用的缓存（constant cache）
• 支持广播——同一个warp内的线程访问同一地址时，只需要一次内存事务

适用场景：存储不会改变的参数，比如卷积核的权重、查找表等。

__constant__ float filter[1024];  // 声明常量内存

__global__ void apply_filter(float* input, float* output) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    // 所有线程访问同一个filter元素时，效率极高
    output[idx] = input[idx] * filter[0] + input[idx+1] * filter[1];
}

纹理内存（Texture Memory）

纹理内存最初是为图形学设计的，但在通用计算中也有独特价值。它提供了：

• 硬件插值（适合图像缩放）
• 边界处理（clamp、wrap等模式）
• 空间局部性优化（2D局部性比全局内存好）

不过说实话，在现代GPU架构中（Volta及以后），纹理内存的优势在逐渐减弱。因为全局内存的缓存机制越来越智能了。我个人现在只在需要硬件插值时才用纹理内存。

内存层次结构总览

下面这张图是我自己画的GPU内存层次结构，你可以直观地看到每一层的位置和关系：

内存访问模式对性能的影响

理解了内存层次结构，我们来看看实际中怎么用。我总结了几个关键原则：

1. 合并访问是王道：全局内存访问一定要合并，这是性能的第一道门槛。
2. 数据重用用共享内存：如果同一个数据要被多个线程使用，先搬到共享内存。
3. 避免寄存器溢出：控制每个线程的变量数量，保持occupancy。
4. 常量数据用常量内存：所有线程访问同一只读数据时，常量内存效率最高。
5. 注意bank conflict：共享内存有32个bank，同时访问同一bank的不同地址会导致冲突。

最后说一句：GPU编程的本质，就是管理数据流动。算力再强，数据喂不进去也是白搭。希望今天的内容能帮你建立起这个思维框架。

想要更多资料，联系可可