1、共享内存基础：GPU内存层次结构、共享内存与全局内存的区别、共享内存的物理位置与延迟特性

各位同学，咱们今天聊聊共享内存。嗯，这是整个优化课程的地基，地基打不牢，后面再花哨的技巧都是空中楼阁。

我个人习惯，讲任何优化之前，先搞清楚“数据在哪”。你想想看，GPU里那么多计算单元，它们怎么拿到数据的？走哪条路？这条路有多快？搞明白这些，你才能写出高效的CUDA代码。

1.1 GPU内存层次结构：一个“金字塔”模型

GPU的内存不是铁板一块，它是个层次结构。我经常跟团队里的新人说，你把GPU想象成一个大型工厂：

• 全局内存（Global Memory）：相当于工厂的中央仓库。容量大（几十GB），但离生产线远，取个零件要跑老远。延迟高，大概几百个时钟周期。
• 共享内存（Shared Memory）：相当于生产线旁边的工具台。每个线程块（Block）独享一个工具台，大家随手就能拿到。延迟低，大概一二十个时钟周期。
• 寄存器（Registers）：相当于工人手里的工具。最快，零延迟（一个周期），但数量极少。
• L1/L2缓存：相当于工厂里的传送带和临时货架。自动管理，但程序员控制不了。

这个金字塔，越往上越快、越贵、越小。咱们优化的核心思路，说白了就是：把最常用的数据，尽量往金字塔上层放。

1.2 共享内存 vs 全局内存：到底差在哪？

我刚开始学CUDA时，觉得这俩不就是内存嘛，能差多少？结果一跑性能数据，傻眼了。咱们直接看对比：

特性	全局内存	共享内存
物理位置	板载DRAM（显存颗粒）	芯片内部SRAM（在SM上）
延迟	~400-800 时钟周期	~20-30 时钟周期
带宽	高（但受限于总线）	极高（每个SM内部）
容量	GB级别	KB级别（典型48KB-164KB）
作用域	所有线程 + Host	单个线程块内
生命周期	程序运行期间	线程块生命周期内
访问模式	需要合并访问（Coalescing）	无合并要求，但存在Bank Conflict

看到没？延迟差了20倍以上。我在项目中遇到过这样一个场景：一个矩阵转置的kernel，直接用全局内存读写，带宽利用率只有20%左右。后来改成用共享内存做中转，带宽直接飙到80%以上。嗯，这就是差距。

1.3 共享内存的物理位置与延迟特性

共享内存到底长什么样？它其实是一块SRAM，直接集成在流多处理器（SM）内部。你想想看，全局内存要走PCIe总线、走显存控制器，而共享内存就在计算单元旁边，物理距离近得可怜。

为什么会这么快？因为SRAM本身速度就比DRAM快，再加上它跟ALU（算术逻辑单元）在同一个芯片上，走的是片上互联，没有那些乱七八糟的总线协议开销。

我记得有一次调试一个性能问题，发现kernel的延迟特别高。用nvprof一分析，发现全局内存访问占了90%的时间。我当时就意识到：这数据根本不应该放在全局内存里。后来我把热点数据搬到了共享内存，延迟从800周期降到了30周期，整个kernel快了将近10倍。

另外，共享内存还有一个隐藏特性：它被划分成32个Bank。同一时刻，每个Bank只能服务一个线程。如果多个线程访问同一个Bank的不同地址，就会发生Bank Conflict，导致访问串行化。这个咱们后面章节会详细讲，今天先记住：共享内存快，但也不是随便用的。

1.4 一个小例子：感受共享内存的威力

光说不练假把式。咱们看一个最简单的例子：向量求和。假设我们要把两个数组a和b对应元素相加，结果存到c里。

纯全局内存版本：

__global__ void vecAdd_global(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

这个版本，每个线程只读一次a和b，写一次c。没有数据复用，用共享内存反而多此一举。但如果是矩阵乘法、卷积、归约这类操作，同一个数据会被多个线程反复读取，共享内存的价值就体现出来了。

共享内存版本（示意）：

__global__ void vecAdd_shared(float *a, float *b, float *c, int n) {
    __shared__ float s_a[256];
    __shared__ float s_b[256];
    
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        s_a[threadIdx.x] = a[idx];
        s_b[threadIdx.x] = b[idx];
        __syncthreads();  // 确保所有线程都加载完毕
        
        c[idx] = s_a[threadIdx.x] + s_b[threadIdx.x];
    }
}

注意看，这里我用了__syncthreads()。这个函数是共享内存的“交通警察”，它保证所有线程都完成加载后，才能进行后续计算。我刚开始写CUDA时经常忘加这个，结果数据还没加载完就开始算，算出来的全是垃圾值。嗯，血的教训。

好了，这一章的基础就讲到这里。下一章咱们深入聊聊共享内存的Bank Conflict，那才是真正考验优化功底的地方。