3. 共享内存详解：Bank Conflict的产生机制、解决策略、动态与静态共享内存分配

共享内存，说白了就是GPU里的一块“片上高速缓存”。它比全局内存快得多，延迟低一个数量级。但用好它，你得先搞懂它的脾气。

我个人习惯把共享内存比作一个“多通道的银行系统”。每个银行（Bank）同一时间只能处理一笔存取款。如果你和你的邻居同时去同一个银行柜台办业务，那就得排队——这就是Bank Conflict的根源。

3.1 Bank Conflict的产生机制

先看硬件结构。共享内存被划分成32个Bank（对应Warp的32个线程）。每个Bank的宽度是4字节（32位）。

举个例子：

// 假设共享内存声明为
__shared__ float shared[32];

// 线程0访问 shared[0]
// 线程1访问 shared[1]
// ...
// 线程31访问 shared[31]

这种情况下，每个线程访问不同的Bank，完美并行。没有冲突。

但如果你这样写：

// 线程0访问 shared[0]
// 线程1访问 shared[0]  // 和线程0抢同一个Bank！
// 线程2访问 shared[0]  // 又抢同一个Bank！

这就产生了Bank Conflict。所有访问同一Bank的线程会被串行化。32个线程抢一个Bank，性能直接降到1/32。

我在项目中遇到过最典型的场景：矩阵转置。直接按列读取、按行写入，Bank Conflict能把性能打掉70%。

3.2 解决策略：实战技巧

解决Bank Conflict，说白了就是“错开访问”。我总结了三种常用手法：

3.2.1 填充法（Padding）

这是最常用的技巧。在声明共享内存时，多申请几个元素，人为制造偏移。

// 错误写法：32列，每行刚好占满32个Bank
__shared__ float tile[32][32];

// 正确写法：每行多申请1个元素，错开Bank对齐
__shared__ float tile[32][32 + 1];  // 每行33个元素

为什么有效？你想想看，原本第0行的第31列和第1行的第0列都在Bank 0。加了1个元素的偏移后，第1行的第0列就跑到了Bank 1。冲突就这么化解了。

3.2.2 改变访问模式

有时候，换个角度思考问题就能解决。比如矩阵转置：

// 冲突版本：按列读取
float val = tile[threadIdx.x][threadIdx.y];  // 同一Warp内，threadIdx.x相同，threadIdx.y不同
// 这会导致同一Bank被多个线程访问

// 优化版本：先读到寄存器，再写共享内存
float reg = input[threadIdx.x][threadIdx.y];
tile[threadIdx.y][threadIdx.x] = reg;  // 转置存储

嗯，这里要注意：改变访问模式往往需要配合数据布局的调整。不能只改读取不改写入。

3.2.3 利用广播机制

有一个特殊情况：如果Warp内所有线程访问的是同一个地址，硬件会触发广播机制。这不叫Bank Conflict，反而是一种优化。

// 所有线程读同一个位置 -> 广播，无冲突
float val = shared[0];

我曾经踩过一个坑：以为广播和冲突是一回事，结果优化了半天反而变慢了。记住，广播是好事，别把它当冲突处理。

3.3 动态与静态共享内存分配

这块我建议你根据实际场景灵活选择。两种方式各有千秋。

3.3.1 静态分配

编译时确定大小，简单直接：

// 静态分配：大小固定
__shared__ float cache[256];

__global__ void kernel() {
    // 直接用
    cache[threadIdx.x] = ...;
}

优点：代码清晰，编译器能优化。缺点：不够灵活，不同调用需要不同大小时得重新编译。

3.3.2 动态分配

运行时指定大小，需要配合extern关键字：

// 动态分配：大小由调用时指定
__shared__ extern float dynamic_cache[];

__global__ void kernel(int size) {
    // 注意：动态分配的共享内存只能声明为一维数组
    // 需要手动计算偏移
    float* my_part = &dynamic_cache[blockIdx.x * size];
}

调用时：

// 第三个参数指定动态共享内存大小（字节）
kernel<<<grid, block, 256 * sizeof(float)>>>(256);

3.3.3 如何选择？

场景	推荐方式	原因
大小固定，所有调用都一样	静态分配	代码简洁，编译器优化好
大小随输入变化	动态分配	灵活，避免浪费共享内存
需要多个不同大小的共享内存	动态分配 + 手动偏移	静态分配只能声明一个动态数组

我个人习惯：能用静态就用静态。只有在需要根据运行时参数动态调整大小时，才用动态分配。毕竟静态分配的代码可读性更好，调试也方便。

3.4 实战建议

最后，分享几个我在项目中积累的经验：

• 先分析，再优化。用NVIDIA Nsight Compute的Memory Workload分析器，一眼就能看出Bank Conflict率。别靠猜。
• 注意Warp边界。同一个Warp内的线程才会产生冲突。不同Warp之间没关系。
• 共享内存不是越大越好。每个SM的共享内存总量有限（通常48KB-164KB）。用多了会降低Occupancy。
• 对齐很重要。访问4字节对齐的地址效率最高。结构体要注意内存对齐。

嗯，共享内存这块内容不少，但核心就这些。下一章我们聊聊常量内存和纹理内存，这两个在特定场景下能给你带来惊喜。