2、GPU内存层次：全局显存、共享内存、寄存器、常量内存、纹理内存

好，咱们今天聊聊GPU的内存层次。说实话，我刚接触GPU编程那会儿，最头疼的就是这堆内存概念。全局显存、共享内存、寄存器……听着就头大。但后来我发现，搞懂它们之间的区别，就像搞懂你办公桌上的东西——有些放抽屉里随手拿，有些得去文件柜翻，还有些得去隔壁仓库搬。速度完全不一样。

GPU的内存层次，说白了就是一套「速度与容量」的权衡体系。离计算单元越近，越快，但容量越小；离得越远，越慢，但能装的东西越多。咱们一个一个来看。

2.1 全局显存（Global Memory）

全局显存，就是大家常说的「显存」。你在显卡参数里看到的8GB、12GB、24GB，指的就是它。它是GPU里最大的存储空间，所有线程都能访问。

特点：

容量最大（几个GB到几十个GB）
延迟最高（几百个时钟周期）
所有线程可读写
没有缓存一致性保证（除非用特殊操作）

⚠️ 注意： 全局显存虽然大，但慢得离谱。我见过不少新手把临时变量全塞全局显存里，结果程序跑得比蜗牛还慢。记住，全局显存是用来存「输入数据」和「最终结果」的，不是给你当临时草稿纸用的。

我在项目中遇到过一个问题：一个图像处理核函数，每次迭代都要从全局显存读一个像素值。结果呢？性能瓶颈全在访存上，计算单元大部分时间都在干等。后来我把数据搬到了共享内存里，速度直接翻了4倍。

2.2 共享内存（Shared Memory）

共享内存是每个线程块（Block）私有的存储空间。同一个块里的线程都能访问它。它的速度比全局显存快得多——大概快20到40倍。

特点：

容量小（通常几十KB到几百KB）
延迟低（几个时钟周期）
同一个Block内的线程可共享
需要手动管理数据搬入搬出

  💡 核心思路： 共享内存就是你的「工作台」。把全局显存里的原材料搬到工作台上，然后所有线程在这个工作台上干活，最后把成品放回全局显存。这样避免了反复去「仓库」取东西。

嗯，这里要注意：共享内存有bank冲突的问题。简单说，如果多个线程同时访问同一个bank的不同地址，就会串行化。我刚开始写代码时经常踩这个坑，后来养成了一个习惯——访问共享内存时，尽量让相邻线程访问相邻地址。这样基本不会冲突。

// 共享内存声明示例
__shared__ float cache[256];

// 典型的规约操作：相邻线程访问相邻地址
int tid = threadIdx.x;
cache[tid] = input[tid];
__syncthreads();

// 规约操作
for (int stride = 1; stride < blockDim.x; stride *= 2) {
    if (tid % (2 * stride) == 0) {
        cache[tid] += cache[tid + stride];
    }
    __syncthreads();
}

2.3 寄存器（Registers）

寄存器是GPU里最快的内存，没有之一。每个线程私有的，编译器自动分配。你写的局部变量，只要不是数组或者结构体，大概率都会被分配到寄存器里。

特点：

速度最快（0延迟，一个时钟周期内完成）
容量极小（每个线程通常几十到几百个）
线程私有，不可共享
编译器自动管理

我个人习惯是：能用寄存器就别用共享内存，能用共享内存就别用全局显存。但寄存器数量有限，每个线程用太多会导致「寄存器溢出」——数据被挤到局部内存里，速度反而变慢。

🔧 小技巧： 在CUDA里可以用 __launch_bounds__ 来限制每个线程块的最大线程数，从而给编译器更多寄存器分配空间。我曾经在一个矩阵乘法核函数里用了这个，性能提升了15%。

2.4 常量内存（Constant Memory）

常量内存，听名字就知道——存常量的。它只有64KB，但有个牛逼的特性：有专门的缓存，而且所有线程读同一个地址时，速度极快。

特点：

容量小（64KB）
只读
有硬件缓存加速
适合存储不会变的参数（如滤波器系数、查找表）

你想想看，如果所有线程都在读同一个系数，比如卷积核的权重，那用常量内存就特别合适。因为硬件会把这个值广播给所有线程，一次读取搞定。

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

// 主机端赋值
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 设备端使用——所有线程读同一个地址，速度飞快
float val = filter[idx];

我曾经在做一个音频滤波项目时，把滤波器系数放在常量内存里，比放在全局显存快了将近3倍。原因很简单：所有线程都在读同一个系数，常量缓存的广播机制完美发挥了作用。

2.5 纹理内存（Texture Memory）

纹理内存，最初是为图形渲染设计的。它有专门的缓存，支持硬件插值、边界处理等功能。对于有空间局部性的数据访问（比如图像像素），纹理内存表现很好。

特点：

有专门的纹理缓存
支持硬件插值（线性、双线性等）
支持边界处理（clamp、wrap等模式）
适合二维/三维空间局部性访问

  💡 适用场景： 图像处理、体渲染、任何需要「邻居数据」的计算。纹理缓存的局部性优化，比手动管理共享内存要省心得多。

我记得有一次做医学图像重建，数据是三维体素。如果用全局显存直接访问，缓存命中率低得可怜。换成纹理内存后，因为相邻线程访问相邻体素，纹理缓存把附近的数据都预取进来了，性能提升了将近2倍。

2.6 内存层次全景图

下面这张图，是我自己画的GPU内存层次结构。你看一眼就明白了：

2.7 如何选择？我的经验总结

说了这么多，到底怎么选？我给大家一个简单的决策流程：

数据是否所有线程共享且只读？ → 常量内存
数据是否有空间局部性（邻居访问）？ → 纹理内存
数据是否被同一个Block反复使用？ → 共享内存
数据是线程私有的临时变量？ → 寄存器
以上都不满足？ → 全局显存（但尽量合并访问）

🔧 避坑指南： 我曾经在一个项目里，把所有数据都放在全局显存里，然后每个线程反复读取。结果性能惨不忍睹。后来我把热点数据搬到共享内存，把常量参数放到常量内存，性能提升了将近10倍。记住：数据搬得越勤快，程序跑得越快。

好了，这一章的内容就到这里。GPU内存层次是CUDA编程的基石，搞懂了它，后面学优化技巧就会轻松很多。下一章咱们聊聊线程调度和warp，那又是另一个有意思的话题。

📚 课程目录

2、GPU内存层次：全局显存、共享内存、寄存器、常量内存、纹理内存

2.1 全局显存（Global Memory）

2.2 共享内存（Shared Memory）

2.3 寄存器（Registers）

2.4 常量内存（Constant Memory）

2.5 纹理内存（Texture Memory）

2.6 内存层次全景图

2.7 如何选择？我的经验总结