共享内存与Bank Conflict：从原理到实战

共享内存，说白了就是GPU上的一块“手边缓存”。它离计算单元很近，访问速度比全局内存快两个数量级。我刚开始接触CUDA时，总觉得共享内存可有可无——反正全局内存也能用嘛。直到有一次，我写了一个矩阵转置的核函数，性能惨不忍睹。加了共享内存后，速度直接翻了十几倍。嗯，从那以后我再也不敢小看它了。

共享内存的工作原理

共享内存是每个线程块私有的存储空间。同一个块里的线程都能访问它。它的生命周期和线程块绑定——块销毁，数据就没了。

硬件上，共享内存被划分成32个等宽的存储单元，叫做Bank。每个Bank的宽度是4字节（32位）。为什么是32个？因为一个Warp正好32个线程。你想想看，这设计是有讲究的。

共享内存的延迟大约是5-10个时钟周期，而全局内存是400-800个周期。差距有多大，不用我多说了吧。

Bank Conflict的产生

Bank Conflict是怎么来的？我举个例子你就明白了。

假设有32个线程，每个线程访问共享内存中的一个float（4字节）。如果线程0访问地址0，线程1访问地址4，线程2访问地址8……这样每个线程访问不同的Bank，完美并行。

但如果线程0和线程1都访问地址0呢？它们就撞到同一个Bank上了。硬件会把这些访问串行化——先服务线程0，再服务线程1。这就是Bank Conflict。

我曾经在项目里遇到过一个坑：用共享内存做归约求和。代码逻辑完全正确，但性能就是上不去。一分析才发现，每个Warp里的线程都在访问同一个Bank位置。加了Padding之后，性能提升了40%。

Bank Conflict的避免方法

避免Bank Conflict，说白了就是让同一个Warp里的线程访问不同的Bank。这里有几种常用方法：

1. 自然对齐访问

最简单的方法。让线程按顺序访问连续地址。比如：

__shared__ float s_data[256];
int tid = threadIdx.x;
float val = s_data[tid];  // 每个线程访问不同的Bank

这种情况下，线程0访问Bank 0，线程1访问Bank 1……完美对齐，没有冲突。

2. Padding技术

有时候你的数据布局会导致冲突。比如二维数组按列访问时，同一行的元素会落在同一个Bank里。这时候可以加一个“填充”元素：

// 不加Padding：每行32个float，按列访问时冲突严重
__shared__ float s_data[32][32];

// 加Padding：每行33个float，按列访问时错开Bank
__shared__ float s_data[32][33];

加了一个float之后，列方向上的访问就错开了。为什么？因为每行多了一个元素，下一行的起始地址偏移了33个float，而不是32个。这样列索引相同的元素就落在了不同的Bank上。

3. Swizzle模式

Swizzle是一种更高级的技巧。它通过重新排列数据在共享内存中的存储位置，来避免Bank Conflict。

我举个例子。假设你要做矩阵转置，输入矩阵按行存储，输出按列存储。如果不做任何处理，读的时候没冲突，但写的时候会有严重的Bank Conflict。

Swizzle的做法是：在写入共享内存时，把数据“打乱”一下。比如：

// 简单的Swizzle示例
__shared__ float s_data[32][33];
int row = threadIdx.x / 32;
int col = threadIdx.x % 32;

// 写入时做Swizzle
s_data[row][col] = input[row * 32 + col];

// 读取时做逆Swizzle
float val = s_data[col][row];  // 转置访问

这里用了Padding（33列）加上行列交换，实际上就是一种Swizzle。更复杂的Swizzle模式会用到异或运算：

// XOR-based Swizzle
int bank = threadIdx.x & 0x1F;
int new_bank = bank ^ (row & 0x1F);
s_data[new_bank][col] = input[row * 32 + col];

这种模式在FFT和某些卷积实现中很常见。我个人习惯在写通用库时用Swizzle，因为它对调用方透明，不需要用户关心底层布局。

动态共享内存分配

有时候你事先不知道需要多少共享内存。比如卷积核的大小是运行时才确定的。这时候就需要动态分配。

动态共享内存的用法很简单：在核函数参数中用extern __shared__声明，然后在调用时指定大小。

// 核函数声明
__global__ void myKernel() {
    extern __shared__ float s_data[];
    // 使用s_data
}

// 调用时指定共享内存大小
myKernel<<<grid, block, 1024 * sizeof(float)>>>();

注意几个要点：

• 动态共享内存只能声明一个数组，类型不限
• 如果你需要多个数组，可以用指针偏移来模拟
• 动态共享内存的大小在编译期不确定，所以不能直接用sizeof

动态分配的一个实用技巧是：把多个数组塞到同一块共享内存里。比如：

__global__ void dualArrayKernel() {
    extern __shared__ char s_data[];
    float* a = (float*)s_data;
    int* b = (int*)(s_data + N * sizeof(float));
    // 现在a和b共享同一块内存
}

这样做的好处是节省共享内存。但要注意对齐问题——float数组的起始地址必须是4字节对齐的。

实战建议

说了这么多，我总结几条实战经验：

1. 先分析，再优化。别一上来就搞Swizzle。先用NVIDIA Nsight Compute看看有没有Bank Conflict，冲突率是多少。
2. Padding是性价比最高的方案。加一个元素就能解决大部分问题，代码改动也小。
3. 动态共享内存适合参数可变的场景。但别忘了检查设备限制。
4. Bank Conflict不是唯一瓶颈。有时候全局内存访问才是大头。先优化全局内存，再考虑共享内存。

这张图把共享内存的知识体系串起来了。从原理出发，到Bank Conflict的产生，再到三种避免方法，最后是动态分配和实战建议。你可以在实际项目中对照这张图，一步步排查问题。

想要更多资料，联系可可