第一讲：CUDA内存模型全景——层次结构与访问延迟对比

各位同学，咱们今天聊点硬核的。CUDA内存模型，说白了就是GPU世界里的一套“交通规则”。你想想看，一个GPU里有几千个核心同时干活，数据怎么存、怎么取、怎么快，全看这套规则玩得转不转。

我个人习惯，在讲任何优化之前，先把内存模型画在纸上。为什么？因为90%的性能瓶颈，都出在“数据放错了地方”。

1.1 全局内存：最熟悉也最慢的“大仓库”

全局内存，就是GPU里那块最大的显存。你调用 cudaMalloc 分配的内存，全在这儿。

• 容量：最大，通常几个GB到几十个GB
• 延迟：约400-800个时钟周期
• 可见性：所有线程都能读写
• 缓存：有L2缓存，但第一次访问基本是“硬伤”

为什么会这么慢？因为全局内存在芯片外面，要通过PCIe或者NVLink连到显存颗粒上。你想想看，从芯片到外面走一圈，能不慢吗？

1.2 共享内存：程序员手中的“高速缓存”

共享内存，是我最喜欢折腾的地方。它就在SM（流多处理器）内部，延迟只有几个时钟周期。

• 容量：每个SM通常48KB或96KB（可配置）
• 延迟：约5-30个时钟周期
• 可见性：同一个线程块内的所有线程共享
• 特点：程序员手动管理，没有自动缓存

我建议你把共享内存当成“临时工作台”。比如做矩阵乘法时，先把数据从全局内存搬到共享内存，然后让线程们在这个小桌子上快速计算。我在优化一个图像处理算法时，就是用共享内存把全局内存访问次数减少了90%，速度直接翻倍。

1.3 寄存器：最快但最稀缺的资源

寄存器，是每个线程私有的“口袋”。延迟几乎为0，一个时钟周期就能拿到数据。

• 容量：每个线程最多255个32位寄存器（实际受限于SM资源）
• 延迟：1个时钟周期
• 可见性：仅当前线程可见
• 特点：编译器自动分配，但你可以用 __launch_bounds__ 控制

这里有个坑：寄存器用多了，会降低“占用率”（occupancy）。什么叫占用率？就是SM上同时活跃的线程数。寄存器用太多，SM能容纳的线程块就少了，反而可能变慢。我曾经为了省几个寄存器，把代码重构了三遍，最后发现性能提升了20%。

1.4 本地内存：寄存器的“溢出区”

本地内存，名字叫“本地”，其实物理上就是全局内存。当寄存器不够用时，编译器会把一些变量“溢出”到本地内存。

• 容量：每个线程最多512KB
• 延迟：和全局内存一样，400-800个时钟周期
• 可见性：仅当前线程可见
• 特点：有L1和L2缓存，但本质还是慢

怎么避免？少用局部大数组，或者用 -maxrregcount 编译选项控制寄存器数量。我见过一个新手写的代码，里面定义了一个 float temp[1024]，结果全被放到了本地内存，性能惨不忍睹。

1.5 常量内存：只读的“广播站”

常量内存，适合存放那些所有线程都要读、但不会改的数据。比如神经网络里的权重参数。

• 容量：64KB
• 延迟：第一次访问约400个时钟周期，之后有缓存
• 可见性：所有线程只读
• 特点：有专用的常量缓存，同一个warp内的线程访问同一地址时，只需一次读取

我建议你把那些“查表”用的数据放常量内存。比如颜色映射表、三角函数表。我在做光线追踪时，把材质参数放常量内存，比放全局内存快了3倍。

1.6 纹理内存：为空间局部性而生

纹理内存，最初是为图形学设计的。它有一些“黑科技”：硬件插值、边界处理、缓存优化。

• 容量：和全局内存共享，但通过纹理缓存访问
• 延迟：有缓存时约100-200个时钟周期
• 可见性：所有线程只读
• 特点：对二维/三维空间局部性访问特别友好

什么时候用纹理内存？当你的线程访问内存地址时，有很强的“空间局部性”。比如图像处理中的卷积操作，每个线程访问一个像素及其邻居。纹理内存的缓存机制，会让这些邻居数据提前加载好。

1.7 实战建议：怎么选？

好了，理论说完了，咱们来点实际的。你在写CUDA代码时，怎么决定数据放哪儿？

1. 先问自己：这个数据是所有线程都要读吗？是 -> 考虑常量内存或纹理内存
2. 再问：同一个线程块内的线程需要共享数据吗？是 -> 用共享内存
3. 再问：数据量小且每个线程独有？用寄存器
4. 最后：实在没地方放，才用全局内存

我个人的经验是：能用共享内存解决的问题，绝不用全局内存。哪怕多写几行代码，也值得。你想想看，400个时钟周期和5个时钟周期的差距，那可是两个数量级。

嗯，这一讲就到这儿。记住一句话：数据放对了地方，性能就成功了一半。下一讲，咱们聊聊怎么用 cudaMalloc 和 cudaMemcpy 玩转全局内存，以及那些你容易忽略的坑。