共享内存优化：Bank Conflict的产生与消除、Padding技术、动态共享内存分配

各位同学，今天我们来聊聊共享内存。说实话，共享内存是CUDA编程里最值得花时间琢磨的东西之一。我见过太多人，核函数写得挺漂亮，结果一跑性能惨不忍睹——十有八九是共享内存没用好。

共享内存为什么重要？因为它快。比全局内存快两个数量级。但快是有代价的——你得会用它。用不好，反而比全局内存还慢。嗯，这里面的门道，我今天一次性讲清楚。

共享内存的硬件结构

先说说共享内存长什么样。每个SM（流多处理器）内部有一块SRAM，这就是共享内存。它被划分成32个bank，每个bank宽度4字节。为什么是32个？因为一个warp正好32个线程。

你想想看，32个线程同时访问共享内存，如果每个线程访问不同的bank，那一次就能搞定。但如果两个线程访问同一个bank的不同地址，那就得串行处理了——这就是bank conflict。

Bank Conflict的产生

我在项目中遇到过这样一个场景：一个简单的矩阵转置，用共享内存做中间缓存。代码逻辑完全正确，但性能只有理论值的30%。一分析，发现bank conflict率高达80%。

为什么会这样？来看个例子：

// 这段代码有严重的bank conflict
__shared__ float s_data[32][32];

int tx = threadIdx.x;
int ty = threadIdx.y;

// 按列写入
s_data[ty][tx] = input[ty * width + tx];

// 按行读取
output[ty * width + tx] = s_data[tx][ty];  // 这里出问题了

问题出在最后一行。当线程按行读取时，s_data[tx][ty]中，同一行的32个元素分布在不同的bank上吗？不是的。因为s_data[tx][ty]的地址计算是 tx * 32 + ty，而bank编号是地址除以4再模32。算下来，同一warp内的线程访问的bank编号是 (tx * 32 + ty) / 4 % 32，当tx变化时，这个值变化很大——但关键是，如果多个线程的tx不同而ty相同，它们会访问同一个bank。

说白了，就是访问模式跟bank的映射关系不匹配。

Bank Conflict的消除

怎么解决？两个思路：

1. 改变数据布局——让连续线程访问连续地址
2. 使用padding——在bank之间插入空位

先说第一种。对于上面的转置问题，我们可以用两次读取来避免conflict：

// 无bank conflict的版本
__shared__ float s_data[32][32];

int tx = threadIdx.x;
int ty = threadIdx.y;

// 按列写入（有conflict，但不可避免）
s_data[ty][tx] = input[ty * width + tx];

__syncthreads();

// 按行读取时，用两次读取避免conflict
float val = s_data[ty][tx];  // 先读一次
output[ty * width + tx] = s_data[tx][ty];  // 再读一次

等等，这看起来没区别啊？嗯，这里的关键是——上面的代码其实还是有conflict。真正有效的做法是使用padding。

Padding技术

Padding说白了就是在数组的每一行末尾加几个空元素，让bank的映射关系错开。我习惯加1个元素就够了：

// 使用padding消除bank conflict
__shared__ float s_data[32][33];  // 注意：33不是32

int tx = threadIdx.x;
int ty = threadIdx.y;

s_data[ty][tx] = input[ty * width + tx];

__syncthreads();

// 现在读取时，地址计算变成 tx * 33 + ty
// bank编号 = (tx * 33 + ty) / 4 % 32
// 因为33不是32的倍数，bank映射被打散了
output[ty * width + tx] = s_data[tx][ty];

我曾经在一个图像处理项目里，用padding把bank conflict从70%降到了5%以下，整体性能提升了2.3倍。就改了一行代码——把[32][32]改成[32][33]。

动态共享内存分配

有时候，共享内存的大小在编译时不确定。比如你要处理不同分辨率的图像，或者不同大小的矩阵。这时候就需要动态分配。

动态共享内存的用法很简单：

// 核函数声明时，使用extern关键字
__global__ void dynamic_shared_kernel(float* input, float* output, int size) {
    // 动态共享内存，大小由第三个启动参数决定
    extern __shared__ float s_data[];
    
    int tid = threadIdx.x;
    if (tid < size) {
        s_data[tid] = input[tid];
    }
    
    __syncthreads();
    
    // 使用共享内存...
    output[tid] = s_data[tid] * 2.0f;
}

// 启动核函数时，指定共享内存大小
int shared_mem_size = blockDim.x * sizeof(float);
dynamic_shared_kernel<<<grid, block, shared_mem_size>>>(d_input, d_output, N);

动态分配的好处是灵活。我做一个自适应滤波算法时，不同图像块的尺寸不一样，用动态共享内存就省去了为每种尺寸写一个核函数的麻烦。

知识体系总览

下面这张图总结了共享内存优化的核心逻辑：

实际项目中的避坑指南

最后，分享几个我踩过的坑：

• 别忘加__syncthreads()——共享内存没有自动同步，不加的话，一个线程写的数据另一个线程可能读不到。我曾经因为这个bug调了整整两天。
• Padding不是万能的——对于某些访问模式，padding反而会降低性能。比如当你的数组宽度本身就是32的倍数时，加padding会浪费带宽。
• 动态共享内存的大小要算清楚——别忘了乘上sizeof(数据类型)。我见过有人传了100，以为是100个float，结果是100字节，只够25个float。
• Bank conflict分析工具要用起来——NVIDIA的Nsight Compute可以直观地看到conflict率。别靠猜，用数据说话。

好了，共享内存这块就讲到这里。记住一句话：共享内存用得好，性能翻倍不是梦；用不好，那就是个昂贵的全局内存缓存。