3. 共享内存与Bank Conflict：性能的隐形杀手
共享内存，说白了就是GPU里的一块“片上高速缓存”。它比全局内存快得多，延迟低一个数量级。但用好它，你得懂它的脾气——尤其是Bank Conflict这个坑。
我个人习惯把共享内存比作一个“多通道仓库”。仓库有32个货架（Bank），每个货架同一时间只能服务一个工人（线程）。如果两个工人同时去同一个货架拿东西，就得排队。这就是Bank Conflict。
3.1 共享内存的Bank组织方式
共享内存被划分成32个Bank（在计算能力5.x及以上的设备上）。每个Bank的宽度是4字节（32位）。相邻的4字节数据存放在相邻的Bank中。
举个例子：
- 地址0-3 → Bank 0
- 地址4-7 → Bank 1
- 地址8-11 → Bank 2
- ... 以此类推
- 地址124-127 → Bank 31
- 地址128-131 → 又回到Bank 0
嗯，这里要注意：地址是按4字节对齐的。如果你用float类型（4字节），那每个float正好占一个Bank的一个槽位。
核心概念：一个Warp（32个线程）访问共享内存时，如果多个线程访问同一个Bank的不同地址，就会发生Bank Conflict。如果所有线程访问同一个地址，那叫Broadcast，不冲突。
我在项目中遇到过一个问题：一个Warp里的线程按列访问一个二维数组。结果32个线程同时访问了同一个Bank，性能直接掉到1/32。当时排查了好久才发现。
3.2 Bank Conflict的产生与消除
Bank Conflict的产生场景其实很典型：
- 步长访问：比如每个线程访问间隔为n的元素。如果n是32的倍数，那所有线程都挤在同一个Bank上。
- 列访问：二维数组按列读取时，如果行数是32的倍数，那每一列的数据都在同一个Bank里。
- 索引计算错误：比如用tid * 32作为偏移量，那所有线程都访问Bank 0。
怎么消除？说白了就是让线程访问的地址均匀分布在不同的Bank上。
常见消除方法：
- 改变访问模式：把列访问改成行访问，或者用转置的方式。
- 使用Padding：在数组每行末尾加几个空元素，让列访问错开Bank。
- 重新组织数据：用共享内存做一次转置，然后再访问。
我的经验：消除Bank Conflict最直接的方法就是加Padding。虽然浪费一点点空间，但换来的是满带宽访问，值！
3.3 Padding技巧
Padding，说白了就是在数组的每一行末尾加几个没用的元素。目的是让列方向上的元素落在不同的Bank里。
举个例子：一个32x32的float数组。如果不加Padding，按列访问时，第0列的所有元素都在Bank 0。加了1个float的Padding后，每行变成33个float。第0列的元素就分布在Bank 0、Bank 1、Bank 2...上了。
// 不加Padding
__shared__ float data[32][32];
// 访问data[tid][0] → 所有线程访问Bank 0 → 32路冲突
// 加Padding
__shared__ float data[32][32 + 1]; // 每行多一个float
// 访问data[tid][0] → 线程0访问Bank 0，线程1访问Bank 1... → 无冲突我曾经在一个矩阵转置的核函数里，就因为没加Padding，性能只有理论带宽的1/10。加了1个float的Padding后，直接跑到80%以上。你想想看，就多了一个float，效果天差地别。
注意：Padding不是越多越好。加太多会浪费共享内存空间，而且可能引入新的对齐问题。一般加1-2个元素就够了。
3.4 动态共享内存分配
动态共享内存，就是运行时才决定共享内存的大小。这在处理不同尺寸的数据时特别有用。
用法很简单：在核函数参数里用 extern __shared__ 声明，然后在调用时指定大小。
// 核函数声明
__global__ void myKernel() {
extern __shared__ float sharedData[];
// 使用sharedData...
}
// 调用时指定大小
int sharedMemSize = blockDim.x * sizeof(float);
myKernel<<<grid, block, sharedMemSize>>>();这里有个坑：动态共享内存只能声明一个数组。如果你需要多个不同类型的数组，得自己手动计算偏移量。
__global__ void myKernel() {
extern __shared__ char sharedMem[];
float* floatData = (float*)sharedMem;
int* intData = (int*)&sharedMem[blockDim.x * sizeof(float)];
// 现在floatData和intData都指向共享内存的不同区域
}我个人习惯用宏或者constexpr来管理偏移量，这样代码更清晰，不容易出错。
避坑指南：动态共享内存的大小是在<<<>>>里指定的，单位是字节。别忘了乘以sizeof(类型)。我曾经因为忘了乘sizeof(float)，结果数据全乱了，排查了半天。
知识体系结构图
这张图把共享内存的核心知识点串起来了。从Bank的组织方式出发，引出Bank Conflict的产生原因，然后给出两种解决方案：Padding和动态分配。你想想看，掌握了这些，是不是对共享内存的优化就有底了？
最后说一句：共享内存优化是CUDA性能调优的重头戏。我见过太多项目，明明算法没问题，就因为共享内存访问模式不对，性能差了好几倍。记住：Bank Conflict是性能的隐形杀手，而Padding是性价比最高的解药。