第四章：共享内存优化——Bank冲突原理、避免策略与动态分配

共享内存，说白了就是GPU里的一块“高速缓存”，由程序员手动控制。它的速度比全局内存快两个数量级，但用不好反而会拖慢性能。我刚开始接触CUDA时，就吃过共享内存的亏——明明数据放进了共享内存，结果比直接用全局内存还慢。后来一查，原来是Bank冲突在作祟。

4.1 Bank冲突原理

共享内存被划分成32个Bank（对应32个线程束线程）。每个Bank的宽度是4字节（32位）。你可以把Bank想象成32个独立的存储通道。理想情况下，每个线程访问不同的Bank，数据就能并行读取。但如果有两个线程访问同一个Bank的不同地址，就会发生冲突。

为什么会这样？因为硬件设计上，每个时钟周期只能从一个Bank读取一个地址。当多个线程同时访问同一个Bank时，访问就会串行化。我遇到过最极端的情况：32个线程全部访问同一个Bank，性能直接降到1/32。

举个例子：

// 无冲突访问：每个线程访问不同的Bank
__shared__ float data[32];
int tid = threadIdx.x;
float val = data[tid];  // 线程0访问Bank0，线程1访问Bank1...

// 2路冲突：两个线程访问同一个Bank
__shared__ float data[64];
int tid = threadIdx.x;
float val = data[tid * 2];  // 线程0和线程32都访问Bank0

嗯，这里要注意：Bank索引的计算方式是 (地址 / 4) % 32。所以访问步长是1时无冲突，步长是2时可能产生2路冲突，步长是32时就是32路冲突——最坏情况。

4.2 避免Bank冲突的策略

我个人习惯用三种方法来避免Bank冲突。每种方法都有适用场景，别死记硬背，要理解背后的原理。

策略一：调整数据布局

最简单的办法就是改变数据在共享内存中的存储方式。比如矩阵转置时，按行存储会导致冲突，按列存储就能避免。

// 冲突版本：按行存储
__shared__ float tile[32][32];
float val = tile[threadIdx.y][threadIdx.x];  // 同一行的线程访问不同Bank，无冲突

// 冲突版本：按列存储
float val = tile[threadIdx.x][threadIdx.y];  // 同一列的线程访问同一个Bank，32路冲突！

// 修复：填充一行
__shared__ float tile[32][33];  // 多出一列
float val = tile[threadIdx.x][threadIdx.y];  // 现在每个线程访问不同的Bank

我在项目中遇到过类似问题：做图像卷积时，按列读取滤波器系数，结果性能惨不忍睹。加了一列填充后，速度提升了3倍。

策略二：使用广播机制

如果所有线程都读取同一个地址，硬件会广播这个值到所有线程，不会产生冲突。这听起来有点反直觉——明明都访问同一个Bank，为什么没冲突？因为硬件专门为这种情况做了优化。

__shared__ float shared_val;
if (threadIdx.x == 0) {
    shared_val = compute_something();
}
__syncthreads();
float val = shared_val;  // 所有线程读取同一个地址，广播模式，无冲突

策略三：改变访问模式

有时候调整线程的索引映射关系就能解决问题。比如把线程ID重新排列，让连续线程访问不连续的Bank。

// 冲突版本
int idx = threadIdx.x * 2;
float val = data[idx];  // 步长为2，2路冲突

// 无冲突版本
int idx = threadIdx.x;
float val = data[idx] + data[idx + 32];  // 每个线程访问两个不同的Bank

你想想看，其实Bank冲突的本质就是“多个线程抢同一个通道”。只要让每个线程的访问地址错开，问题就解决了。

4.3 动态共享内存分配

动态共享内存，说白了就是在运行时才确定共享内存的大小。静态分配是在编译时写死的，不够灵活。我建议在以下场景使用动态分配：

• 输入数据大小在编译时未知
• 同一个内核需要处理不同大小的数据块
• 想要提高代码的通用性

使用方法很简单：

// 内核声明时使用 extern
__global__ void dynamic_shared_kernel() {
    extern __shared__ float shared_data[];
    // 现在shared_data的大小由启动配置决定
}

// 启动内核时指定共享内存大小
int shared_size = block_size * sizeof(float);
dynamic_shared_kernel<<<grid, block, shared_size>>>();

动态分配时要注意对齐问题。共享内存的起始地址是16字节对齐的，但如果你在动态共享内存中存储多种数据类型，需要手动处理对齐。

// 在动态共享内存中存储多种类型
extern __shared__ char shared_mem[];
float* float_data = (float*)shared_mem;
int* int_data = (int*)(shared_mem + block_size * sizeof(float));
// 注意：int_data可能没有4字节对齐，需要手动调整

我曾经在项目中用动态共享内存实现了一个通用的矩阵乘法内核，可以处理任意大小的矩阵分块。虽然代码复杂了些，但维护起来特别方便。

4.4 性能调优实战建议

最后，分享几个我多年积累的实战经验：

1. 先用工具检测：使用 nvcc --ptxas-options=-v 查看共享内存使用量，用Nsight Compute分析Bank冲突。
2. 不要过度优化：如果共享内存访问不是瓶颈，花太多精力消除Bank冲突反而得不偿失。我见过有人为了消除冲突，把代码改得面目全非，结果性能提升不到5%。
3. 注意寄存器压力：有时候消除Bank冲突会增加寄存器使用量，导致占用率下降。要权衡利弊。
4. 测试不同配置：共享内存大小、块大小、线程布局，这些参数相互影响。我习惯写一个自动化测试脚本，跑几十组参数找最优解。

嗯，这一章的内容就到这里。下一章我们会讲全局内存的合并访问，那是另一个容易踩坑的地方。到时候再聊。