显存基础：GPU显存架构概览

做GPU优化这么多年，我最大的感触就是——不了解显存，就别谈性能。很多同学一上来就盯着算法复杂度，结果代码跑起来比CPU还慢，问题往往出在数据搬运上。

GPU的显存体系，说白了就是一个金字塔结构。越往上，速度越快，容量越小；越往下，容量越大，延迟越高。我刚开始接触CUDA时，总觉得显存就是显存，哪来这么多门道？直到有一次，我写了个矩阵乘法，全局内存读写毫无优化，结果带宽利用率不到10%。嗯，那次教训挺深刻的。

  核心原则： 尽量让数据靠近计算单元，减少对全局内存的访问。

全局内存（Global Memory）

全局内存是GPU里最大的存储区域，所有线程都能访问。但它的延迟很高——大约400-800个时钟周期。你想想看，CPU的L1缓存访问只要几个周期，这差距有多大。

我个人习惯把全局内存想象成一个大仓库。数据从CPU传过来，先放仓库里。线程要干活，得从仓库搬数据。搬一次很慢，所以最好一次搬一大块。

这里有个关键概念——合并访问。什么意思？就是相邻的线程访问相邻的地址。我见过太多新手代码，线程0访问地址0，线程1访问地址100，线程2访问地址200...这种写法，带宽直接腰斩。

// 好的写法：合并访问
__global__ void good_kernel(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[idx];  // 相邻线程访问相邻地址
    // ... 处理数据
}

// 差的写法：非合并访问
__global__ void bad_kernel(float* data) {
    int idx = threadIdx.x * 100 + blockIdx.x;  // 跨度100
    float val = data[idx];  // 带宽利用率极低
    // ... 处理数据
}

我的经验： 写全局内存访问时，先画个线程-地址映射图。保证线程ID和地址偏移量是线性关系，基本就对了。

共享内存（Shared Memory）

共享内存是GPU里的高速缓存，由程序员手动管理。它在同一个block内的线程间共享，延迟只有几十个周期，比全局内存快一个数量级。

我记得有个项目做图像滤波，原始写法每次从全局内存读一个像素的邻域数据，性能惨不忍睹。后来改成用共享内存缓存整个tile，速度提升了8倍。为什么？因为数据复用了。

共享内存的使用套路其实很固定：

从全局内存加载数据到共享内存（需要同步）
在共享内存上做计算
写回全局内存

__global__ void shared_mem_kernel(float* input, float* output, int N) {
    __shared__ float tile[256];  // 声明共享内存
    
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        tile[threadIdx.x] = input[idx];  // 加载到共享内存
    }
    __syncthreads();  // 同步！确保所有线程都加载完毕
    
    // 在共享内存上做计算
    float result = tile[threadIdx.x] * 2.0f;
    
    output[idx] = result;
}

注意： 共享内存容量有限，每个SM通常只有几十KB到一百多KB。我曾经在项目里贪心，想缓存更多数据，结果共享内存用超了，kernel启动失败。教训是：先算好容量，再动手写代码。

寄存器（Registers）

寄存器是GPU里最快的存储，没有之一。每个线程独享，延迟为0（实际上就是指令操作数）。但数量极其有限——每个线程最多255个寄存器，实际可用通常更少。

我见过一个常见的坑：变量声明太多，导致寄存器溢出（spilling）。溢出后，变量会被存到本地内存（实际在全局内存里），速度骤降。怎么判断？看编译报告里的"register spilling"警告。

避坑指南： 我曾经在写一个复杂的卷积核时，用了大量临时变量。编译后发现性能比预期差很多，一查是寄存器溢出。解决办法：把一些中间结果合并，或者用共享内存分担。

本地内存（Local Memory）

本地内存听起来像"本地"的，其实它物理上就在全局内存里。只是每个线程私有的。当寄存器不够用时，编译器会自动把变量放到本地内存。

为什么会用到本地内存？常见情况：

数组索引不是编译期常量
结构体太大，寄存器放不下
递归函数调用

嗯，这里要注意：本地内存的延迟和全局内存一样高。所以尽量避免。我一般会检查编译器的ptx输出，看看有没有"local"关键字。

常量内存（Constant Memory）

常量内存用于存储只读数据，所有线程都能访问。它有个特殊机制——当所有线程访问同一个地址时，会被广播，速度极快。但如果每个线程访问不同地址，性能反而比全局内存还差。

我常用的场景：存储卷积核参数、查找表、归一化系数等。这些数据在整个kernel执行期间不变，而且所有线程读同一个值。

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

__global__ void apply_filter(float* input, float* output) {
    int idx = threadIdx.x;
    float val = input[idx];
    float result = val * filter[0];  // 所有线程读同一个地址
    output[idx] = result;
}

  关键点： 常量内存只有64KB，别想存太多东西。而且更新常量内存需要调用cudaMemcpyToSymbol，不能在kernel里修改。

纹理内存（Texture Memory）

纹理内存最初是为图形渲染设计的，但用在通用计算里也有奇效。它支持硬件插值和边界处理，特别适合图像处理、插值计算等场景。

我记得有个项目做图像缩放，用纹理内存的硬件双线性插值，比自己写插值函数快了3倍。而且纹理内存对空间局部性的访问模式有优化——相邻线程访问相邻像素时，缓存命中率很高。

不过纹理内存也有局限：只能读不能写，而且API稍微复杂些。我个人习惯是：只有需要插值或边界处理时，才考虑纹理内存。普通数据访问，共享内存更灵活。

知识体系总览

下面这张图，是我自己总结的GPU显存架构。每次做优化前，我都会对照着看一遍，确认数据该放哪。

总结一下

GPU显存优化，说白了就是让数据待在它该待的地方。频繁访问的数据放寄存器或共享内存，只读且所有线程访问相同地址的放常量内存，需要插值的放纹理内存。全局内存？尽量少碰。

我刚开始做优化时，总想着用各种奇技淫巧。后来发现，把基础打好，把数据放对位置，性能自然就上来了。你想想看，如果数据搬运占了90%的时间，算法再优化也没用，对吧？

我的建议： 拿到一个kernel，先分析数据访问模式。哪些数据是复用的？哪些是只读的？哪些是线程私有的？画个数据流图，再决定用哪种内存。这一步花10分钟，可能省下10小时的调试时间。

想要更多资料，联系可可

📚 课程目录

显存基础：GPU显存架构概览

全局内存（Global Memory）

共享内存（Shared Memory）

寄存器（Registers）

本地内存（Local Memory）

常量内存（Constant Memory）

纹理内存（Texture Memory）

知识体系总览

总结一下