一、CUDA内存模型全景：从全局内存到寄存器，一张图看懂GPU存储层次

做GPU开发这些年，我见过太多人一上来就写核函数，结果性能惨不忍睹。问他们为什么慢，十个有九个答不上来。其实说白了，CUDA编程的核心就两个字——内存。你搞懂了GPU怎么存取数据，性能优化就成功了一大半。

今天咱们先从整体上看看CUDA的内存模型长什么样。别急着写代码，先把这张地图刻在脑子里。

1.1 GPU存储层次：一个金字塔

GPU的内存不是铁板一块，它分了好多层。我习惯把它想象成一个金字塔：

• 顶层：寄存器（最快，但极小）
• 中间层：共享内存、L1/L2缓存
• 底层：全局内存（最慢，但最大）

为什么会这样分层？说白了就是成本和速度的博弈。寄存器快到飞起，但一个SM里就那么几百个。全局内存容量大，但延迟是寄存器的上百倍。

嗯，这里要注意：不同层次的内存，对程序员的可控性完全不同。寄存器你基本管不了，编译器帮你搞定。共享内存你得手动管理。全局内存嘛，你得操心怎么把数据搬进去、搬出来。

1.2 一张图看懂存储层次

我画了张SVG图，把整个存储层次串起来。你看完这张图，再往下读会轻松很多。

这张图我画了好几个版本，最后选了金字塔形式。你想想看，越往上走，速度越快，但容量越小。这就是GPU存储层次的核心矛盾。

1.3 各层内存速览

咱们从上到下，一个一个过一遍。我挑重点说，不啰嗦。

寄存器（Registers）

这是离计算单元最近的地方。每个线程独享，编译器自动分配。我在项目中遇到过一个问题：寄存器溢出。就是线程用了太多寄存器，编译器被迫把数据 spill 到本地内存，性能直接腰斩。所以写核函数时，变量别声明太多，能用共享内存就别死抱着寄存器不放。

共享内存（Shared Memory）

同一个线程块内的线程可以共享。这是手动优化的主战场。我曾经在一个矩阵转置的项目里，用共享内存把带宽利用率从30%拉到了85%。说白了，共享内存就是程序员手里的瑞士军刀——用好了事半功倍，用不好反而拖慢速度。

共享内存分两种：静态分配和动态分配。静态的用 __shared__ 声明，动态的靠核函数参数传进去。我个人习惯用静态的，代码更清晰。

// 静态共享内存
__shared__ float s_data[256];

// 动态共享内存（核函数调用时指定大小）
extern __shared__ float s_data_dyn[];

L1/L2缓存

这两个是硬件自动管理的，程序员只能间接影响。比如你可以通过 cudaFuncSetAttribute 调整L1和共享内存的比例。嗯，这里要注意：L1缓存对全局内存的访问有加速效果，但对共享内存没用——因为共享内存本身就是手动管理的缓存。

全局内存（Global Memory）

这是GPU里最大的存储空间，也是性能瓶颈的重灾区。所有线程都能访问，但延迟高达几百个周期。怎么优化？核心就两条：合并访问和减少访问次数。

合并访问什么意思？就是让相邻的线程访问相邻的地址。GPU一次取128字节，如果你让线程0取地址0，线程1取地址4，线程2取地址8……那一次内存事务就搞定了。反之，如果线程0取地址0，线程1取地址1000，那就要两次事务，性能直接打对折。

常量内存和纹理内存

这两个是特殊用途的。常量内存适合存只读的、所有线程都访问相同地址的数据（比如系数表）。纹理内存适合二维空间局部性强的访问模式（比如图像处理）。

我个人用得不多，但知道它们的存在很重要。遇到合适的场景，它们能带来意外之喜。

1.4 各层内存对比表

我把关键参数整理成了一张表，方便你对照着看：

内存类型	作用域	生命周期	延迟（cycles）	典型容量	缓存
寄存器	单个线程	线程内	1	255个/线程	无
共享内存	线程块	块内	~30	48-164 KB/SM	无（手动管理）
L1缓存	SM	自动	~30	与共享内存共享	硬件管理
L2缓存	芯片	自动	~200	几MB	硬件管理
全局内存	所有线程+主机	程序运行期间	400-800	GB级	L1/L2
常量内存	所有线程+主机	程序运行期间	~200（缓存命中）	64 KB	专用常量缓存
纹理内存	所有线程+主机	程序运行期间	~200（缓存命中）	取决于设备	专用纹理缓存

你看这张表，延迟从1到800，差了三个数量级。所以写CUDA程序时，脑子里要时刻想着：我当前的数据在哪一层？能不能往上挪一挪？

1.5 一个简单的例子

光说不练假把式。咱们看一个最简单的向量加法，感受一下内存层次的影响：

// 全局内存版本（慢）
__global__ void vecAdd_global(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];  // 每次都要访问全局内存
    }
}

// 共享内存版本（快）
__global__ void vecAdd_shared(float *a, float *b, float *c, int n) {
    __shared__ float s_a[256], s_b[256];
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int tid = threadIdx.x;
    
    // 一次加载到共享内存
    if (idx < n) {
        s_a[tid] = a[idx];
        s_b[tid] = b[idx];
    }
    __syncthreads();
    
    // 从共享内存读取，快得多
    if (idx < n) {
        c[idx] = s_a[tid] + s_b[tid];
    }
}

第二个版本虽然多了 __syncthreads() 同步开销，但后续的加法操作直接从共享内存读数据，延迟从400 cycles降到了30 cycles。数据量越大，优势越明显。

1.6 小结

好了，这一章咱们把CUDA内存模型的全景看了一遍。从寄存器到全局内存，每一层都有它的脾气。记住金字塔结构，记住延迟差异，记住合并访问原则——这些是后续所有优化的基础。

下一章我会深入全局内存，聊聊怎么写出带宽利用率90%以上的代码。嗯，到时候见。