3、共享内存与Bank冲突：共享内存架构、Bank映射机制、冲突避免策略（Padding、Swizzle）

共享内存，说白了就是GPU里的一块“片上高速缓存”。它离计算单元特别近，延迟比全局内存低一个数量级。我刚开始接触CUDA时，觉得这玩意儿不就是个快一点的存储器吗？后来踩了坑才明白——共享内存用得好不好，性能差距能到10倍以上。

嗯，今天咱们就聊聊共享内存的核心机制，以及那个让无数开发者头疼的Bank冲突问题。

3.1 共享内存的物理架构

先看一张图（脑补一下）：共享内存被划分成32个Bank，每个Bank宽度是4字节（32位）。为什么是32个？因为一个Warp正好32个线程。你想想看，这设计是有意为之的——让每个线程访问不同的Bank，就能并行读写。

每个Bank本质上是一个单端口SRAM。这意味着什么？同一时钟周期内，一个Bank只能响应一次访问请求。如果两个线程同时访问同一个Bank，那就得串行化处理，这就是Bank冲突的根源。

关键点：共享内存的带宽是32 × 4字节 = 128字节/周期。如果没有冲突，一个Warp的32个线程可以同时完成数据访问。但一旦发生冲突，带宽就会打折。

我在项目中遇到过一种情况：一个简单的矩阵转置，用共享内存做中间缓冲，结果性能比直接读全局内存还慢。排查了半天，发现就是Bank冲突导致的——每个线程都在访问同一个Bank的不同地址，冲突率高达100%。

3.2 Bank映射机制

Bank的映射规则其实很简单：地址除以4（因为每个Bank存4字节），然后对32取模。

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

举个例子：

  • 地址0x00 → Bank 0
  • 地址0x04 → Bank 1
  • 地址0x08 → Bank 2
  • ...
  • 地址0x7C → Bank 31
  • 地址0x80 → Bank 0（又回来了）

为什么会这样？因为地址是连续映射的。这种设计对顺序访问非常友好——线程0读地址0，线程1读地址4，线程2读地址8...完美并行。

但问题来了：如果每个线程访问的数据间隔是4字节的整数倍，且步长是32的倍数呢？比如每个线程读一个float（4字节），但步长是128字节（32个float）。那所有线程都会访问同一个Bank！

注意：Bank冲突只发生在同一个Warp内部。不同Warp之间不存在Bank冲突，因为它们访问共享内存的时间是错开的。

3.3 冲突的几种模式

我习惯把Bank冲突分成三类：

冲突类型 描述 性能影响
无冲突 每个线程访问不同的Bank 100%带宽
n路冲突 n个线程访问同一个Bank 带宽降至1/n
广播 所有线程读同一个地址 无冲突（硬件支持广播）

这里有个容易忽略的点：广播模式。如果Warp内所有线程都读同一个地址，硬件会把这个地址广播到所有线程，不产生冲突。但写操作不行——多个线程写同一个地址，结果未定义。

3.4 冲突避免策略

好了，理论说完了，咱们聊聊实战。怎么避免Bank冲突？我总结了三种常用方法。

3.4.1 Padding（填充）

这是最直观的方法。既然冲突是因为地址映射到同一个Bank，那我们在数组里加一些“无用”的填充数据，把地址错开不就行了？

举个例子：声明共享内存时，本来需要 32 × 32 的float数组，我们声明成 32 × 33。

// 有冲突的版本
__shared__ float tile[32][32];

// 无冲突的版本（加Padding）
__shared__ float tile[32][33];  // 每行多一个float

为什么这样有效？原本每行32个float，正好占满32个Bank。下一行的第0列又映射到Bank 0。加了1个float的填充后，下一行的第0列映射到Bank 1，完美错开。

我的经验：Padding的代价是浪费了一点共享内存。但32×33只比32×32多了3%的存储，换来的是100%的带宽利用率，这笔账怎么算都划算。

3.4.2 Swizzle（交错映射）

Padding虽然好用，但有时候我们不想浪费那点空间。这时候可以用Swizzle——通过改变数据的存储布局，让原本会冲突的访问模式变成无冲突。

Swizzle的核心思想：把数据重新排列，使得原本映射到同一个Bank的数据分散到不同Bank。

// 一个简单的Swizzle示例
// 假设我们要访问一个32×32的矩阵，按列访问
// 原始布局：tile[row][col] → Bank = col % 32
// Swizzle后：tile[row][(col + row) % 32] → Bank = (col + row) % 32

// 写入时做Swizzle
__shared__ float swizzled_tile[32][32];
int swizzled_col = (col + row) & 31;  // 等价于 % 32
swizzled_tile[row][swizzled_col] = data;

// 读取时做逆Swizzle
int original_col = (swizzled_col - row) & 31;
float val = swizzled_tile[row][swizzled_col];

你看，通过一个简单的异或或加法操作，就把列访问的冲突变成了无冲突。代价是多了几次整数运算——但相比Bank冲突带来的性能损失，这点计算开销微乎其微。

我曾经在一个图像处理项目里用过Swizzle。那个算法需要按列访问共享内存中的像素数据，Padding会浪费L1缓存空间。用Swizzle后，性能提升了40%，而且没有额外存储开销。

3.4.3 改变数据访问模式

有时候，最简单的办法是改变线程的访问模式。比如：

  • 把列访问改成行访问
  • 让线程处理连续的数据块
  • 使用向量化加载（如 float4）

举个例子：

// 有冲突的列访问
float val = shared_mem[threadIdx.x][threadIdx.y];  // 按列读

// 改成行访问（无冲突）
float val = shared_mem[threadIdx.y][threadIdx.x];  // 按行读

当然，这取决于算法本身是否允许你改变数据布局。有些算法天生就需要按列访问，那就只能用Padding或Swizzle了。

3.5 实战避坑指南

最后，分享几个我踩过的坑：

我曾经...在一个项目里用了共享内存做归约（Reduction），结果发现性能上不去。排查后发现，归约的每一步都在访问同一个Bank——因为所有活跃线程都在读同一个索引。解决方案很简单：让每个Warp内的线程访问不同的Bank，或者用Warp Shuffle指令代替共享内存。

另一个常见坑：Bank冲突的检测。NVIDIA的profiler（如Nsight Compute）可以显示Bank冲突的次数。我建议你在优化时一定要跑一下profiler，别靠猜。

还有一点：共享内存的大小限制。每个SM的共享内存是有限的（比如48KB或96KB）。Padding和Swizzle都会增加共享内存的使用量，别一不小心用超了。

我的习惯：先写一个无优化的版本，跑profiler看Bank冲突率。如果冲突率超过10%，就考虑用Padding或Swizzle。如果冲突率很低（比如小于1%），那就不值得优化——毕竟，过早优化是万恶之源。

好了，共享内存和Bank冲突就聊到这儿。记住一句话：共享内存是GPU性能的倍增器，Bank冲突是它的阿喀琉斯之踵。理解了映射机制，掌握了Padding和Swizzle，你就能在大多数场景下写出高效的共享内存代码。

← 上一章 📖 返回目录 下一章 →