返回首页共享内存 被模拟的访问
Bank Conflict 模拟器
输入访问模式与 padding，看到每个 bank 上落了几个线程、最大冲突路数以及需要多少个周期。
共享内存访问模式
模式
32
0
经典的 +1 让行宽变成奇数，与 32 互质
最大冲突路数
32×
访问要串行 32 个周期
用到的 bank
1 / 32
存在 32 路冲突，这次共享内存访问需要 32 个周期而不是 1 个，有效带宽降到 3%。 试试把 padding 调成 1 ，让行宽变成奇数。
32 个 bank 的落点
每一列是一个 bank，柱子上的每个方块是一个访问它的线程。柱子越高，串行周期越多。
0
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
0
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
每个线程的映射
t0
w0
bank 0
t1
w32
bank 0
t2
w64
bank 0
t3
w96
bank 0
t4
w128
bank 0
t5
w160
bank 0
t6
w192
bank 0
t7
w224
bank 0
t8
w256
bank 0
t9
w288
bank 0
t10
w320
bank 0
t11
w352
bank 0
t12
w384
bank 0
t13
w416
bank 0
t14
w448
bank 0
t15
w480
bank 0
t16
w512
bank 0
t17
w544
bank 0
t18
w576
bank 0
t19
w608
bank 0
t20
w640
bank 0
t21
w672
bank 0
t22
w704
bank 0
t23
w736
bank 0
t24
w768
bank 0
t25
w800
bank 0
t26
w832
bank 0
t27
w864
bank 0
t28
w896
bank 0
t29
w928
bank 0
t30
w960
bank 0
t31
w992
bank 0
CUDA C++
1__shared__ float tile[32][32];2 3float v = tile[threadIdx.x][0]; // 按列访问4// word = tid * 32 → bank = (tid * 32) % 325// 行宽 32 与 32 的最大公约数 > 1 → 仍有冲突模型说明：这里按 4 字节 bank 宽度、32 个 bank 计算，对应 float 类型的共享内存访问。64 位的 double 访问在现代架构上由硬件拆成两个阶段处理，规则略有不同，实践中同样通过 padding 或 swizzle 解决。