第三章 内存模型精讲：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存的层次结构与访问特性

各位同学，今天我们来聊聊CUDA的内存模型。说实话，我见过太多新手在内存使用上栽跟头。我自己刚入行时也踩过不少坑——有一次一个核函数跑得比CPU还慢，查了两天才发现是全局内存访问模式不对。嗯，内存用好了，性能翻倍；用不好，GPU再强也白搭。

3.1 内存层次结构概览

CUDA的内存不是铁板一块，它分了好几个层级。你想想看，GPU上有几千个核心同时干活，如果大家都去抢同一块内存，那不乱套了？所以NVIDIA设计了一套分层的内存体系，从快到慢、从近到远，各有各的用途。

我个人习惯把CUDA内存想象成一个金字塔：

• 寄存器：最快，每个线程私有，容量极小
• 共享内存：次快，同一个线程块内共享，可编程控制
• 本地内存：寄存器不够用时自动使用，速度慢
• 全局内存：所有线程都能访问，容量最大，但延迟最高
• 常量内存：只读，有缓存加速，适合广播数据
• 纹理内存：专门为空间局部性优化，有硬件插值

下面这张图是我自己画的，把整个内存层次串起来了：

3.2 全局内存——最常用也最坑

全局内存是GPU上的"大仓库"，所有线程都能读写。但它的延迟高达400-800个时钟周期。我在项目中遇到过，有人把全局内存当寄存器用，结果性能惨不忍睹。

访问模式是关键。全局内存最喜欢"合并访问"——也就是同一个warp里的32个线程，访问连续的内存地址。举个例子：

// 好的写法：合并访问
float data = global_array[threadIdx.x];  // 连续地址

// 坏的写法：非合并访问
float data = global_array[threadIdx.x * 32];  // 跨步访问

3.3 共享内存——可编程的"小快灵"

共享内存是每个线程块私有的高速缓存，延迟只有5-10个周期。说白了，它就是让你手动管理的一块"L1缓存"。容量很小，通常只有几十KB，但用好了能解决大问题。

我建议你在以下场景优先考虑共享内存：

• 线程块内需要数据共享
• 全局内存访问模式不连续
• 需要做归约、扫描等操作

共享内存有个"银行冲突"的问题，得小心。举个例子：

// 共享内存声明
__shared__ float s_data[256];

// 无银行冲突：每个线程访问不同bank
float val = s_data[threadIdx.x];

// 有银行冲突：两个线程访问同一bank的不同地址
float val = s_data[threadIdx.x * 2];  // 可能冲突

3.4 寄存器——最快的"贴身侍卫"

寄存器是每个线程私有的，访问延迟只有1个周期。但数量有限——每个线程最多255个寄存器（具体看计算能力）。如果你声明太多局部变量，编译器会把多余的"溢出"到本地内存，速度一下就掉下来了。

我记得有一次优化一个核函数，发现占用率只有25%。查了一下，每个线程用了128个寄存器。后来手动拆分成多个小核函数，寄存器降到32个，占用率升到100%，吞吐量翻了三倍。

3.5 本地内存——寄存器的"备胎"

本地内存听起来像"本地"，其实它物理上就在全局内存里。只是每个线程有自己的地址空间，看起来像私有的。什么时候会用？

• 寄存器不够用，编译器自动溢出
• 声明了大型局部数组（比如 float arr[100]）
• 使用了取地址操作（&var）

说白了，本地内存就是"伪装的全局内存"。能用寄存器就别用它。

3.6 常量内存——广播利器

常量内存只有64KB，但有个好处：有专门的缓存，而且同一个warp里的线程读同一个地址时，只需要一次内存访问。适合存放查找表、系数矩阵这类只读数据。

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

// 在主机端初始化
cudaMemcpyToSymbol(coeff, host_coeff, sizeof(float) * 256);

// 核函数中使用
float val = coeff[index];  // 如果warp内index相同，一次读取搞定

3.7 纹理内存——为空间局部性而生

纹理内存最初是为图形学设计的，但用在通用计算上也有奇效。它有几个特点：

• 有硬件缓存，对二维/三维空间局部性友好
• 支持硬件插值（线性插值、双线性插值）
• 边界处理方便（clamp、wrap等模式）

我在做图像处理项目时，用纹理内存做双线性插值，比手动实现快了4倍。不过要注意，纹理内存是只读的，而且绑定和解绑有开销。

3.8 各内存对比一览

内存类型	位置	延迟	容量	作用域	可读写
寄存器	SM内部	1周期	~256KB/SM	线程私有	读写
共享内存	SM内部	5周期	~48-164KB/SM	线程块	读写
本地内存	全局内存	~400周期	~512KB/线程	线程私有	读写
全局内存	显存	~400周期	~80GB	所有线程	读写
常量内存	显存+缓存	~5周期(缓存命中)	64KB	所有线程	只读
纹理内存	显存+缓存	~5周期(缓存命中)	~80GB	所有线程	只读

3.9 实战建议

说了这么多，总结几条我自己的经验：

1. 先想清楚数据流：写代码前，画个图，标出数据从哪里来、到哪里去、中间谁需要共享。
2. 全局内存访问要合并：这是性能的底线，做不到合并，其他优化都是白搭。
3. 共享内存当缓存用：把频繁访问的全局数据先搬到共享内存，再让线程块内反复使用。
4. 寄存器别贪多：每个线程超过64个寄存器就要警惕了，看看能不能拆函数或者减少局部变量。
5. 常量内存适合小数据广播：64KB虽然小，但放个查找表绰绰有余。
6. 纹理内存适合空间局部性强的场景：图像、矩阵、网格数据，纹理内存有奇效。