Kernel 3: Bank Conflict Free

通过填充消除共享内存 bank 冲突

共享内存 Bank 详解

内存组织

GPU 共享内存分为 32 个 bank（现代架构）。每个 bank 每时钟周期可服务一次访问。

地址 → Bank 索引:  address % 32

Bank 0  Bank 1  ...  Bank 31
┌─────┐ ┌─────┐     ┌─────┐
│ [0] │ │ [1] │ ... │ [31]│  ← 地址 0-31
├─────┤ ├─────┤     ├─────┤
│ [32]│ │ [33]│ ... │ [63]│  ← 地址 32-63
├─────┤ ├─────┤     ├─────┤
│ ... │ │ ... │ ... │ ... │
└─────┘ └─────┘     └─────┘

冲突场景

cpp

__shared__ float tile[32][32];

// 在内积循环中：
for (int k = 0; k < 32; ++k) {
    sum += tile[ty][k] * tile[k][tx];  // 所有线程访问第 k 列
}

当 warp 中的线程读取 tile[k][0], tile[k][1], ..., tile[k][31] 时：

线程 0 访问地址：k * 32 + 0 → Bank (k * 32) % 32 = 0
线程 1 访问地址：k * 32 + 1 → Bank (k * 32) % 32 = 0
...
线程 31 访问地址：k * 32 + 31 → Bank (k * 32) % 32 = 0

结果：所有 32 个线程同时命中 Bank 0 → 32 路冲突！

性能影响

指标	Tiled (32×32)	Bank-Free (32×33)	改进
GFLOPS (1024³)	753	673	略有变化
Bank 冲突	32 路	无	已消除
共享内存	8 KB	8.4 KB	+5.5% 开销
访问周期	32×	1×	32× 更快

Bank-free kernel 在不同场景下提供更一致的性能。

关键要点

32 个 Bank：共享内存分为 32 个 bank（现代 GPU）
冲突：多个线程命中同一 bank 时访问串行化
填充：第二维加 1 将步长从 32 变为 33
公式：Bank 索引 = (row × (TILE_SIZE + 1) + col) % 32
开销：仅 3% 更多共享内存换取 32× 性能提升

Kernel 3: Bank Conflict Free ​

共享内存 Bank 详解 ​

内存组织 ​

冲突场景 ​

性能影响 ​