__shared__ 数组比 global memory 快很多——延迟跟 L1 cache 同级(~20 cycle)。但它有个几乎所有 CUDA 新手都会踩的陷阱:bank conflict。
跟 coalescing 一起,这是 CUDA 性能两大支柱。
1. 为什么 shared memory 要分 bank
shared memory 物理上是 SRAM,单口读写一次只能服务一个 lane。但 warp 内 32 个 lane 通常同一周期同时发出 shared memory 请求——硬件怎么处理?
答案:把 shared memory 切成 32 个独立的 bank,每个 bank 自己有一个端口。32 个 lane 同时访问 32 个不同 bank → 全部并行,1 cycle 完成。
Shared memory(假设 128 字节):
地址: 0 1 2 ... 31 32 33 ... 63 64 65 ...
│ │ │ │ │ │ │ │ │
▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼
bk0 bk1 bk2 ... bk31 bk0 bk1 ... bk31 bk0 bk1 ...32 个 bank 循环映射——每 4 字节切换一个 bank,32 个 bank 用完之后回到 bank 0。
2. Bank 映射规则(必须背)
对 32-bit(4 字节)元素:
bank_id = (address_in_bytes / 4) % 32
= element_index % 32 (如果元素是 float / int 这种 4 字节)举例:__shared__ float A[256];
| 元素索引 | 字节地址 | bank |
|---|---|---|
A[0] | 0 | 0 |
A[1] | 4 | 1 |
A[31] | 124 | 31 |
A[32] | 128 | 0(回头) |
A[33] | 132 | 1 |
A[64] | 256 | 0 |
两个 A[i] 和 A[j] 落在同一 bank ⟺ (i - j) % 32 == 0(对 4 字节元素)。
3. Bank conflict 是什么
warp 内多个 lane 在同一周期访问同一个 bank 的不同地址 → 硬件无法并行,串行化为 N 次访问。
只有同一 bank 的不同地址才冲突。两个特殊情况不冲突:
| 情况 | 为什么不冲突 |
|---|---|
| 同 bank 同地址(广播) | 硬件支持 broadcast,1 次读多个 lane 收到结果 |
| 不同 bank | 各自独立通道,完全并行 |
冲突程度:
| 含义 | 延迟 | |
|---|---|---|
| No conflict | 32 lane → 32 个不同 bank,或全广播 | 1× |
| 2-way conflict | 32 lane 中 2 个落同 bank | 2× |
| N-way conflict | 最坏:32 lane 全落同一 bank | 32× |
4. 典型冲突场景
4.1 跨度访问
__shared__ float A[1024];
int tid = threadIdx.x;
float v = A[tid * 2]; // ← stride 2A[0 × 2] = A[0] → bank 0
A[1 × 2] = A[2] → bank 2
A[16 × 2] = A[32] → bank 0(又是 0!)
A[17 × 2] = A[34] → bank 2(又是 2!)
warp 内 32 个 lane 落进 16 个 bank,每个 bank 2 个 lane,2-way conflict。
通用规律:stride = 2^k(k ≥ 1)→ 2^k-way conflict。stride = 32 是最惨,32-way conflict(全堆在 bank 0)。
4.2 矩阵转置(经典)
__shared__ float tile[32][32];
tile[ty][tx] = global_in[...]; // 写:tile[ty][0..31] → 各 bank,无冲突
__syncthreads();
float v = tile[tx][ty]; // 读:tile[0..31][ty] → 全是 bank ty,32-way conflict!写的时候按行访问 → 每 lane 不同 bank → 无冲突。
读的时候按列访问 → 32 个 lane 全访问列 ty,字节地址都是 ty * 4 + i * 128,全落 bank ty,32 倍变慢。
4.3 不冲突的反例
__shared__ float A[1024];
float v = A[tid]; // stride 1,32 lane 落 32 个不同 bank,无冲突
float v = A[tid + 32]; // 同上(整体偏移,bank 也偏移,还是 32 个不同 bank)
float v = A[0]; // 全部广播,无冲突5. 解决方案
5.1 Padding(最常用)
__shared__ float tile[32][33]; // ← 33 不是 32!
tile[ty][tx] = ...; // 写:仍然无冲突
float v = tile[tx][ty]; // 读:每行多了 1 个元素,列访问的字节地址错开了 33*4 = 132
// 132 % 128 = 4,bank 错开[TILE][TILE + 1] 是 CUDA 里地标级的 idiom。多一列内存,换全列访问无冲突。
5.2 Swizzle(更高级,无 padding 浪费)
把存储位置按 xor 模式打乱:
int swizzled_col = col ^ (row & (TILE - 1));
tile[row][swizzled_col] = val;
// 读取时用同样的 xor 算回不浪费 padding 空间,但需要算 xor 索引。GEMM / Conv 这种对 shared memory 用量敏感的 kernel 常用。
5.3 改变访问模式
如果是算法允许,把”按列遍历”改成”按行遍历”。bank conflict 通常只在某一种访问方向上发生,翻一下方向就没了。
6. 怎么验证
用 Nsight Compute:
| 指标 | 找哪 | 含义 |
|---|---|---|
Shared Load Bank Conflicts | Memory Workload Analysis | 总冲突数 |
Shared Store Bank Conflicts | 同上 | 写冲突 |
Bank Conflict ÷ Shared Load | 算比例 | 看占比 |
| Source 标签 | 点对应代码行 | 直接定位哪条 __shared__ 访问产生冲突 |
跑:
ncu --section MemoryWorkloadAnalysis ./your_app理想值:0 conflicts。如果有但很少(< 几个百分比),可能不值得花时间消;大量冲突(尤其 16/32 way)→ 必须修。
7. 8-byte 元素的特殊情况
double / int64 / float2 是 8 字节。这种元素一次访问跨两个 bank:
double 地址 0 → bank 0 + bank 1
double 地址 8 → bank 2 + bank 3合并规则仍然按 bank 算冲突,但单次访问占两个 bank,warp 实际需要 2 周期完成 1 次访问(等价于”硬件层面的”2-way conflict)。这是 fp64 在 shared memory 上比 fp32 慢的内在原因之一。
float4 / int4 等 16 字节元素同理 → 4 周期。
8. 跟 coalescing 的关系
CUDA 性能两大支柱:
| global memory | shared memory | |
|---|---|---|
| 机制 | warp 内 32 lane 的请求合并成 1 次 DRAM 事务 | 32 个 bank 并行,同 bank 不同地址要串行 |
| 优化方向 | 让 thread index 沿最内维度走 | tile 加 padding [TILE][TILE+1] / swizzle |
| 代价 | 跨度访问慢 32 倍 | bank conflict 慢 N 倍(N = 同 bank 冲突 lane 数) |
| 诊断指标 | Sectors/Req、Global Load Throughput | Shared Load Bank Conflicts |
写 GPU kernel 的标准 checklist:
- 数据布局:SoA、向量 load,让 global 访问合并
- tile 排布:
[TILE][TILE+1]或 swizzle,让 shared 访问无 bank conflict
两个都做对了,memory bound kernel 性能基本拿到。
9. 一句话
shared memory 是 32 个并行 bank,bank 编号 =
(地址 / 4) % 32。warp 内多 lane 同时访问同 bank 不同地址 → N-way 串行。最常见的修法是[TILE][TILE+1]加一列 padding 把列访问的 bank 错开。