4、共享内存(Shared Memory)优化：Bank Conflict原理与解决、静态与动态分配、Padding技术、同步原语(__syncthreads)

共享内存，说白了就是GPU里的一块“手边缓存”。它离计算单元近，速度比全局内存快两个数量级。但用好它，你得懂它的脾气。

我刚开始做CUDA优化时，总觉得把数据搬进共享内存就万事大吉了。结果一跑性能分析，发现带宽利用率只有30%。嗯，问题就出在Bank Conflict上。

4.1 共享内存的硬件结构：Bank到底是什么？

共享内存在硬件上被划分成32个等宽的存储模块，每个模块叫一个Bank。你可以把Bank想象成32个独立的小仓库，每个仓库一次只能服务一个线程。

为什么是32个？因为一个Warp正好32个线程。设计者的意图很明确：让Warp里的每个线程都能同时访问不同的Bank，实现全带宽。

每个Bank的宽度是4字节（32位）。也就是说，地址0-3属于Bank 0，地址4-7属于Bank 1，依此类推。计算公式很简单：Bank索引 = (地址 / 4) % 32。

4.2 Bank Conflict：性能杀手

当一个Warp里的多个线程访问同一个Bank的不同地址时，就发生了Bank Conflict。硬件会把冲突的访问串行化，一次只能处理一个。冲突越多，性能越差。

举个例子，假设线程0访问地址0，线程1访问地址32。地址0在Bank 0，地址32也在Bank 0（因为32/4=8，8%32=8？等等，我算错了）。

重新算：地址0 → Bank 0，地址32 → 32/4=8，8%32=8，所以是Bank 8。没冲突。

那什么情况会冲突？线程0访问地址0，线程1访问地址4。地址0在Bank 0，地址4在Bank 1。也没冲突。

真正冲突的是：线程0访问地址0，线程1访问地址128。128/4=32，32%32=0，都在Bank 0。这就冲突了。

访问模式	Bank冲突数	有效带宽利用率
无冲突（每个线程访问不同Bank）	0	100%
2路冲突	2	50%
4路冲突	4	25%
32路冲突（最坏情况）	32	3.125%

4.3 常见冲突模式与解决方法

4.3.1 行主序访问 vs 列主序访问

二维数组的访问模式最容易引发冲突。假设我们有一个32×32的float数组，每个线程处理一行：

// 行主序访问：无冲突
__shared__ float data[32][32];
int tid = threadIdx.x;
float val = data[tid][0];  // 线程0访问Bank 0，线程1访问Bank 1...完美

但如果改成列主序：

// 列主序访问：32路冲突！
__shared__ float data[32][32];
int tid = threadIdx.x;
float val = data[0][tid];  // 所有线程都访问Bank 0！

为什么会这样？因为data[0][tid]的地址是&data[0] + tid * sizeof(float)。tid从0到31，地址间隔4字节，正好落在Bank 0到Bank 31。等等，不对。

我重新算一下：data[0][0]在Bank 0，data[0][1]在Bank 1，data[0][2]在Bank 2……data[0][31]在Bank 31。这其实没冲突！

真正有冲突的是data[0][tid]当数组是float data[32][32]时，每行32个float占128字节，正好是32个Bank。所以data[0][tid]和data[1][tid]的Bank索引是一样的。如果Warp里线程访问不同行的同一列，就会冲突。

4.3.2 步长访问

步长访问也很常见。比如处理矩阵转置时，步长等于矩阵宽度。如果宽度是32的倍数，就会产生严重冲突。

// 步长为32：32路冲突
__shared__ float tile[32][32];
int idx = threadIdx.x + blockIdx.x * blockDim.x;
int idy = threadIdx.y + blockIdx.y * blockDim.y;
float val = tile[threadIdx.y][threadIdx.x];  // 转置访问

解决方法还是Padding：

// 加一列Padding，消除冲突
__shared__ float tile[32][33];  // 注意第二维是33
int idx = threadIdx.x + blockIdx.x * blockDim.x;
int idy = threadIdx.y + blockIdx.y * blockDim.y;
float val = tile[threadIdx.y][threadIdx.x];  // 现在没冲突了

4.4 Padding技术详解

Padding说白了就是在数组的每一行末尾加几个没用的元素，让行的长度不再是32的倍数。这样同一列的元素就分散到不同的Bank里了。

Padding的大小怎么选？一般加1个元素就够了。但如果你用的是double（8字节），可能需要加更多。

数据类型	原始行宽	Padding后行宽	Bank冲突消除情况
float (4B)	32	33	完全消除
double (8B)	16	17	完全消除
int (4B)	64	65	完全消除

4.5 静态分配 vs 动态分配

共享内存的分配方式有两种：静态和动态。我个人习惯用静态分配，因为代码更清晰。但有些场景必须用动态。

4.5.1 静态分配

在核函数里直接用__shared__声明，大小在编译时就确定了：

__global__ void kernel() {
    __shared__ float cache[256];  // 编译时确定大小
    // ...
}

优点是简单直观，缺点是大小固定，不够灵活。

4.5.2 动态分配

在核函数外面用extern __shared__声明，大小在启动核函数时指定：

__global__ void kernel() {
    extern __shared__ float cache[];  // 大小由启动参数决定
    // ...
}

// 启动时指定共享内存大小
kernel<<<grid, block, 256 * sizeof(float)>>>();

动态分配的好处是同一个核函数可以处理不同大小的数据。我在做自适应算法时经常用这个技巧。

4.6 同步原语：__syncthreads()

共享内存是Block内所有线程共享的。这就带来了一个问题：一个线程写，另一个线程读，怎么保证顺序？

__syncthreads()就是干这个的。它是一个屏障，Block里所有线程都必须到达这个点，才能继续往下走。

__global__ void kernel() {
    __shared__ float cache[256];
    
    // 线程0写入
    if (threadIdx.x == 0) {
        cache[0] = 3.14f;
    }
    
    __syncthreads();  // 确保写入完成
    
    // 所有线程都能读到正确的值
    float val = cache[0];
}

我曾经在一个项目里忘了加__syncthreads()，结果有些线程读到的是旧值，有些读到的是新值。调试了两天才找到原因。嗯，从那以后我每次用共享内存都会检查同步点。

另外要注意，__syncthreads()会消耗性能。别在不需要的地方滥用。我见过有人每个语句后面都加一个，结果性能直接腰斩。

4.7 知识体系总览

下面这张图把共享内存优化的核心知识点串起来了：

4.8 实战建议

说了这么多，总结几条我在项目中积累的经验：

• 先分析再优化：用NVIDIA Nsight Compute跑一下，看看共享内存的Bank Conflict计数。别凭感觉猜。
• Padding是银弹：遇到列访问冲突，先试试加一列Padding。90%的情况都能解决。
• 同步点别乱放：只在需要的地方加__syncthreads()。放多了性能下降，放少了数据出错。
• 动态分配要谨慎：动态分配的共享内存不能初始化，而且多个动态数组需要手动计算偏移。我一般只用静态。