第3章共享内存（Shared Memory）详解

3、共享内存（Shared Memory）详解：共享内存的bank结构、bank冲突的产生与避免、共享内存的配置与大小

共享内存，说白了就是GPU里的一块“片上高速缓存”。它离计算单元非常近，延迟极低。我个人习惯把它比作“团队协作的白板”——同一个线程块里的所有线程，都能在上面读写数据。但用不好，性能会大打折扣。

3.1 共享内存的Bank结构

共享内存内部被划分成多个Bank。你可以把每个Bank想象成一个独立的内存通道。每个Bank在同一时钟周期内，只能响应一次访问请求。

以NVIDIA的架构为例，每个Bank的宽度是32位（4字节）。共享内存的总线宽度就是 Bank数量 × 32位。常见的配置有16个Bank或32个Bank。

举个例子，假设共享内存有32个Bank，地址0-3字节映射到Bank0，地址4-7字节映射到Bank1，以此类推。地址128-131字节又回到Bank0。这是典型的低位交叉编址方式。

核心公式：

Bank索引 = (字节地址 / 4) % Bank总数

嗯，这里要注意：地址是按4字节对齐的。如果你访问一个float变量，它天然就对齐到4字节边界。但如果你访问一个char数组，就得小心了。

3.2 Bank冲突的产生与避免

Bank冲突，是共享内存性能的“头号杀手”。

什么是Bank冲突？

当一个Warp（32个线程）里的多个线程，在同一时钟周期内，访问同一个Bank的不同地址时，就会发生Bank冲突。这些访问会被串行化，导致有效带宽下降。

举个例子：

// 假设共享内存 __shared__ float s_data[32];
// 线程0访问 s_data[0]  -> Bank0
// 线程1访问 s_data[1]  -> Bank1
// ...
// 线程31访问 s_data[31] -> Bank31
// 这是完美情况，无冲突

再看一个冲突的例子：

// 线程0访问 s_data[0]  -> Bank0
// 线程1访问 s_data[32] -> Bank0  (因为32%32=0)
// 线程2访问 s_data[64] -> Bank0
// 三个线程都访问Bank0，发生3路Bank冲突

我在项目中遇到过这样一个坑：做矩阵转置时，按列写入共享内存，结果每个线程都访问同一个Bank的不同地址，性能直接掉到原来的1/16。后来改成按行写入，再按列读出，完美避开了冲突。

避坑指南：

我曾经在做一个图像滤波算法时，没注意共享内存的步长。线程0访问s_data[0]，线程1访问s_data[16]，线程2访问s_data[32]... 结果步长16正好是Bank数的约数，导致严重的Bank冲突。后来我加了一个“padding”技巧，在每行末尾多分配一个元素，把步长错开，冲突就消失了。

如何避免Bank冲突？

连续访问：让相邻线程访问相邻地址。这是最理想的情况。
使用padding：在数组每行末尾加几个无用元素，改变地址到Bank的映射关系。
改变数据布局：比如用结构体数组（SoA）代替数组结构体（AoS）。
利用广播机制：如果多个线程访问同一个地址，硬件会广播，不产生冲突。

小技巧：

我建议你在写代码时，先算一下每个线程访问的Bank索引。如果发现多个线程的Bank索引相同，但地址不同，就说明有冲突。可以用一个简单的公式：bank_id = (thread_id * stride) % bank_count。当stride和bank_count互质时，冲突最少。

3.3 共享内存的配置与大小

共享内存的大小，不是你想设多少就设多少。它受硬件限制。

以NVIDIA的架构为例：

架构	每SM最大共享内存	每线程块最大共享内存	Bank数
Kepler (SM 3.x)	48 KB	48 KB	32
Maxwell (SM 5.x)	96 KB	48 KB	32
Pascal (SM 6.x)	96 KB	48 KB	32
Volta/Turing (SM 7.x)	96 KB	48 KB	32
Ampere (SM 8.x)	128 KB	48 KB	32

你想想看，每个SM的共享内存是有限的。如果你一个线程块用了太多共享内存，那SM上能同时运行的线程块数量就会减少。这叫“占用率下降”。

如何配置共享内存大小？

在CUDA中，有两种方式：

静态分配：在核函数里用__shared__关键字声明。大小在编译时就固定了。
动态分配：在调用核函数时，用第三个尖括号参数指定大小。例如：kernel<<<grid, block, shared_mem_size>>>()。

举个例子：

// 静态分配
__global__ void myKernel() {
    __shared__ float s_data[256];  // 固定1KB
}

// 动态分配
__global__ void myKernel() {
    extern __shared__ float s_data[];  // 大小由调用时决定
}

int main() {
    myKernel<<<1, 256, 1024>>>();  // 分配1KB共享内存
}

经验之谈：

我个人习惯在项目初期，先用cudaDeviceGetAttribute查询硬件的共享内存上限。然后根据算法需求，算出一个合理的值。别贪心，够用就行。我曾经为了省一点全局内存访问，把共享内存塞得满满的，结果占用率从100%掉到25%，性能反而更差了。

共享内存与L1缓存的配置

在一些架构上，共享内存和L1缓存共享同一块物理存储。你可以通过cudaFuncSetAttribute或cudaDeviceSetCacheConfig来调整它们的比例。

比如在Pascal架构上，你可以选择：

cudaFuncCachePreferShared：优先给共享内存，比如48KB共享 + 16KB L1
cudaFuncCachePreferL1：优先给L1缓存，比如16KB共享 + 48KB L1
cudaFuncCachePreferEqual：均衡分配，比如32KB共享 + 32KB L1

嗯，这里要注意：不是所有架构都支持动态调整。Volta和之后的架构，共享内存和L1是独立的，不用你操心。

调试建议：

我建议你用cudaOccupancyMaxPotentialBlockSize这个API，它能帮你算出在给定共享内存大小下，能达到的最大占用率。别自己瞎猜，让工具说话。