第4章共享内存优化

4. 共享内存优化：共享内存与全局内存的区别、Bank Conflict、共享内存的Padding技巧

各位同学，今天我们来聊聊共享内存。说实话，在GPU编程里，共享内存是个又爱又恨的东西。爱它，是因为它比全局内存快得多；恨它，是因为用不好反而会拖慢性能。我在调优一个矩阵乘法的时候，就曾经被Bank Conflict坑了整整两天。

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

先说说最基础的区别。全局内存，说白了就是显存，容量大（几个GB到几十个GB），但延迟高（几百个时钟周期）。共享内存呢，是每个线程块内部的一块高速缓存，容量小（通常几十KB），但延迟低（几个时钟周期）。

我习惯这么理解：全局内存就像去图书馆借书，来回跑一趟要花不少时间；共享内存就像把书放在自己的办公桌上，随手就能拿到。你想想看，如果频繁访问同一批数据，是不是应该先搬到桌上再说？

具体差异我整理了个表格：

特性	全局内存	共享内存
延迟	~400-800 cycles	~5-30 cycles
带宽	高（但受限于访存模式）	极高（无合并访问限制）
容量	GB级别	通常16-96 KB per SM
作用域	所有线程	单个线程块内
生命周期	程序运行期间	线程块执行期间

嗯，这里要注意：共享内存虽然快，但容量极其有限。我见过有人把整个数据集都塞进共享内存，结果编译都过不了。所以，共享内存是用来做数据复用的，不是用来做数据存储的。

4.2 Bank Conflict：共享内存的隐形杀手

共享内存之所以快，是因为它被分成了32个Bank（在大多数现代GPU上）。每个Bank可以在一个时钟周期内响应一次访问。理想情况下，32个线程同时访问32个不同的Bank，完美并行。

但问题来了——如果两个线程访问同一个Bank的不同地址，就会发生Bank Conflict。这时候，访问会被串行化，性能直接打折扣。

我曾经在一个图像处理项目里，把共享内存用得很顺手，结果一跑profiling发现，Bank Conflict导致性能下降了40%。当时我那个郁闷啊……

来看个具体的例子：

// 这段代码会导致严重的Bank Conflict
__shared__ float shared_data[32][32];

// 线程 (tx, ty) 访问 shared_data[tx][ty]
// 如果按列访问，同一列的数据都在同一个Bank里
float val = shared_data[threadIdx.x][threadIdx.y];  // 糟糕！

为什么会这样？因为共享内存的地址到Bank的映射是：Bank = (address / 4) % 32。对于二维数组 shared_data[32][32]，shared_data[0][0] 和 shared_data[1][0] 的地址差是 32 * 4 = 128 字节，128 / 4 = 32，正好是Bank数量的整数倍。所以它们落在同一个Bank上！

核心结论：Bank Conflict的本质是多个线程访问了同一个Bank的不同地址。如果所有线程访问同一个地址（广播），反而没有冲突。

4.3 Padding技巧：简单粗暴的解决方案

怎么解决Bank Conflict？我个人的经验是：Padding（填充）是最简单有效的方法。

思路很简单：在数组的每一行末尾加几个无用元素，让地址偏移不再是Bank数量的整数倍。这样，原本落在同一个Bank的地址就被打散了。

还是刚才的例子：

// 加1个元素的Padding，把32列变成33列
__shared__ float shared_data[32][33];  // 注意：33列！

// 现在按列访问，地址偏移变成了 33 * 4 = 132 字节
// 132 / 4 = 33，33 % 32 = 1，完美错开！
float val = shared_data[threadIdx.x][threadIdx.y];  // 没有Bank Conflict了

我建议你记住这个公式：如果数组宽度是Bank数量的整数倍，就加1个元素的Padding。对于32个Bank的情况，就是把宽度从N改成N+1。

小技巧：Padding浪费的空间其实很少。比如32x32的数组，加1列Padding只浪费了3%的空间，但可能带来30%以上的性能提升。这笔账怎么算都划算。

4.4 实战：用Padding优化矩阵转置

咱们来看一个完整的例子。矩阵转置是典型的访存密集型操作，用共享内存做分块转置时，Bank Conflict非常常见。

// 没有Padding的版本
__global__ void transpose_no_padding(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    // 加载到共享内存（按行读）
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 从共享内存写出（按列写）—— 这里会有Bank Conflict！
    out[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

// 有Padding的版本
__global__ void transpose_with_padding(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE + 1];  // 加Padding
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 现在按列写也没有Bank Conflict了
    out[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

我在实际测试中，这个简单的改动让转置操作的吞吐量提升了将近1.5倍。嗯，有时候优化就是这么朴实无华。

4.5 知识体系总览

最后，我用一张图来总结本章的核心逻辑。这张图展示了共享内存优化的完整路径：从理解差异，到识别问题，再到动手解决。

避坑指南：我曾经犯过一个错误——在动态分配的共享内存上做Padding。动态共享内存的地址是运行时确定的，Padding需要手动计算偏移量，很容易搞错。建议优先用静态共享内存，省心。

好了，关于共享内存优化，核心就这些。记住三点：共享内存是数据复用的利器，Bank Conflict是性能的隐形杀手，Padding是最简单的解药。下次写CUDA代码的时候，不妨多留个心眼，看看你的共享内存访问有没有踩坑。

公众号:蓝海资料掘金营，微信deep3321

📚 课程目录

4. 共享内存优化：共享内存与全局内存的区别、Bank Conflict、共享内存的Padding技巧

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

4.2 Bank Conflict：共享内存的隐形杀手

4.3 Padding技巧：简单粗暴的解决方案

4.4 实战：用Padding优化矩阵转置

4.5 知识体系总览