Bank Conflict（存储体冲突 / Bank 冲突）是 GPU 计算（如 NVIDIA CUDA）中共享内存（Shared Memory）访问时极易遇到的另一种典型性能瓶颈。

什么是 Bank Conflict？

为了在多线程并发访问时提供极高的带宽，GPU 的Shared Memory（共享内存）在物理硬件上被划分为若干个大小相等、可独立访问的存储单元，这些单元被称为Bank（存储体）。

在 NVIDIA GPU 架构中，Shared Memory 通常划分为32 个 Bank（刚好对应一个 Warp 中的 32 个线程）。连续的 32-bit 模式 word（4 字节）按顺序循环映射到这 32 个 Bank 中：

• Word 0→\rightarrow→Bank 0
• Word 1→\rightarrow→Bank 1
• …
• Word 31→\rightarrow→Bank 31
• Word 32→\rightarrow→Bank 0（地址地址对 128 字节取模循环）

当一个 Warp（32 个线程）在同一时钟周期内发起 Shared Memory 的读写请求时：

• 无冲突（No Conflict）：若 32 个线程分别访问 32 个不同的 Bank，或者多个线程访问同一个 Bank 中的完全相同的数据地址（触发广播/Multicast 机制），所有的内存请求可以并行完成，达到峰值带宽。
• 发生冲突（Bank Conflict）：若 32 个线程中有 2 个或更多线程访问同一个 Bank 中的不同内存地址，硬件就无法在单一周期内并行处理这些请求。

硬件如何处理 Bank Conflict？

发生 Bank Conflict 时，硬件会将针对同一个 Bank 的不同请求进行串行化（Serialization）：

• 如果有NNN个线程访问了同一个 Bank 的不同地址（称为NNN-way Bank Conflict），该 Bank 的访问过程就会被拆分为NNN个独立的时钟周期依次进行。
• 这会导致共享内存的访问延迟成倍增加，整体吞吐量下降至原来的1/N1/N1/N。

典型代码场景分析

1. 0-way Bank Conflict（完美无冲突）

__shared__ float s_data[1024]; // 线程 i 访问 s_data[i] // 线程 0 -> Bank 0 (s_data[0]) // 线程 1 -> Bank 1 (s_data[1]) // ... // 线程 31 -> Bank 31 (s_data[31]) float val = s_data[threadIdx.x];

结果：32 个线程精准映射到 32 个不同的 Bank，1 个时钟周期完成读写。

2. 2-way / N-way Bank Conflict（步长不当导致冲突）

在矩阵转置或图像处理等跨步长（Stride）访问 Shared Memory 的场景中非常常见：

__shared__ float s_data[1024]; // 访问步长为 2 (stride = 2) // 线程 0 -> s_data[0] -> Bank 0 // 线程 1 -> s_data[2] -> Bank 2 // ... // 线程 16 -> s_data[32] -> Bank 0 (与线程 0 冲突！) // 线程 17 -> s_data[34] -> Bank 2 (与线程 1 冲突！) float val = s_data[threadIdx.x * 2];

结果：发生2-way Bank Conflict。对于步长stride = 32（如按列读取按行存储的32x32矩阵），32 个线程会全部挤在 Bank 0 上，造成极其严重的32-way Bank Conflict。

3. 特例：广播机制（Broadcast）与 无冲突

// 所有线程访问相同的地址 s_data[0] (都在 Bank 0) float val = s_data[0];

结果：虽然所有线程都映射到 Bank 0，但由于访问的是完全相同的地址，硬件会触发广播（Broadcast）机制，一次性将数据发给所有线程，不会发生 Bank Conflict。

常见的优化与规避策略

1. 内存填充（Padding）

这是解决二维 Shared Memory 矩阵访问冲突最经典且成本最低的技术。

假设有一个32 x 32的共享内存数组，按列读取时跨步为 32，会导致 32-way Bank Conflict：

// 原始声明：32 列 __shared__ float tile[32][32]; // tile[i][threadIdx.x] 会导致严重冲突

通过在每一行尾部手动填充 1 个虚拟元素（将列数改为 33）：

// 填充后声明：33 列 __shared__ float tile[32][33];

原理：行宽变成 33 后，下一行的第 0 列元素在物理内存中的索引移动了 1，映射到的 Bank 也顺移了 1 个（Bank = index % 32）。此时按列访问tile[i][col]时，原本映射到相同 Bank 的元素被错开到了不同的 Bank 中，冲突瞬间消除！

2. 改变数据布局与访问模式

• Swizzling（地址错位/交错）：利用异或（XOR）等逻辑运算对索引进行重映射（如在 Tensor Core / Cutlass 库中常用的 Layout Swizzling），使得二维数据的访问逻辑在物理 Bank 上打散。

3. 利用动态 Bank 宽度配置（针对部分架构）

某些 GPU 架构允许配置 Shared Memory 的 Bank 宽度（如从 4-byte 切换到 8-byte）。如果线程是以double或float2（8 字节）为单位访问 Shared Memory，调整 Bank 宽度可以更好地匹配数据类型。