小李第一天到一家 AI 芯片公司报到，就被派去优化一个图像处理内核。同事扔给他一句话：“程序慢，多半是共享内存的'柜台'堵了。”小李一头雾水：什么柜台？同事把他带到 GPU 大楼的“收发室”——共享内存。这里有 32 个柜台，每个柜台一位营业员，每轮能同时服务 32 位顾客（线程）。可是小李很快发现一个怪现象：明明顾客们办的是完全不同的业务，却总有一股脑儿挤在同一个柜台前排长队的时刻，剩下 31 个柜台闲着。老工程师翻开叫号规则，一句话点醒他：柜台号 = 业务号 mod 32。当数据按矩阵的列去取时，所有线程分到的柜台号一模一样——挤爆 0 号柜台，这就是 Bank Conflict。解法也出人意料地简单：在每行末尾加一个“占位隔板”（Padding），让取号结果错开，冲突自动消散。这篇文章就带你把这门“柜台拥堵学”彻底搞懂：从共享内存的组织、冲突的分类与代价，到诊断方法、Padding 优化和访问模式改写，一步步成为 GPU 性能优化团队里合格的大堂经理。

岗位关联

掌握 Bank Conflict 的成因与解法，你就能把 CUDA 内核从“能跑”提升到“跑得快”——这是 AI 训练、推理加速与高性能计算岗位面试中的高频加分硬技能。

目录
  1. 可视化演示
  2. 共享内存基础
  3. Bank组织原理
  4. 冲突类型
  5. 诊断方法
  6. Padding优化
  7. 消除策略
  8. AI 应用场景
  9. Python 演示
  10. 常见误区
  11. 面试题

可视化演示

动态演示Bank Conflict的核心机制，通过基础/核心/应用/扩展四种模式展示关键概念与交互流程。

本节配有交互动效（进入视口后自动加载，约 3KB）。

Bank Conflict：共享内存 共享内存Bank组织结构 Bank Conflict类型 核心机制可视化 Bank Conflict整体架构与关键组件

共享内存基础

GPU 的每个 SM（流式多处理器，可以理解成大楼里的一层办公区）都自带一个“内部收发室”——共享内存。同一楼层不同工位的员工（同一 block 内的线程）要互相传文件，根本不必跑去几十公里外的中央仓库（全局内存），直接在本楼层收发室交接就行。取一个包裹，去收发室只要几秒钟，跑一趟中央仓库要几小时。正因为这个收发室就在“片上”，它的延迟比全局内存低一个数量级以上，成了 GPU 编程中最高效的临时数据交换站。

共享内存的容量并不大，一般只有几十 KB，但它有两个黄金特性：第一，它位于 SM 内部，是典型的“片上高速存储”，访问延迟远低于显存；第二，它对本 block 内的所有线程可见，线程们能通过它互相传话、协作计算。在 CUDA 里，我们常常把全局内存的数据分块搬进共享内存，像收发室把大宗货物先卸到大厅分拣，再让员工就近领取。

那么这部“电梯”到底有多快？硬件给共享内存设计的并行能力是：每个时钟周期，32 个柜台同时各服务一位顾客，每个柜台一次递出 4 字节。因此理论带宽就是 32 乘以 4 字节。

收发室虽然近，但真正决定效率的，是来办事的人会不会把柜台挤爆。小李隐隐觉得，这个“柜台分布”里藏着大文章。

Bshared=32×4 字节=128 字节/周期\displaystyle B_{shared} = 32 \times 4\ \text{字节} = 128\ \text{字节/周期}
这个公式算的是“收发室每秒钟最多能递出多少东西”：32 个柜台同时开工、每个柜台每次递 4 字节，乘积就是硬件承诺的峰值吞吐。
  1. 32 个柜台并行一个时钟周期内，32 个银行同时开工、互不等待，这是共享内存吞吐的根基。
    32 banks\displaystyle 32\ \text{banks}
  2. 每柜台 4 字节每个柜台一次只服务一个“词”，正好是一个 float 或 int 的宽度，对应共享内存的 4 字节访问粒度。
    4 字节/柜台\displaystyle 4\ \text{字节/柜台}
  3. 相乘得带宽32 个 4 字节并行搬运，总带宽 128 字节/周期。看起来很美——前提是柜台别被挤爆。
    128 字节/周期\displaystyle 128\ \text{字节/周期}

Bank组织原理

收发室大厅里，32 个柜台一字排开，每个柜台只有一位营业员，一次只能服务一位顾客的一个“词”（4 字节）。柜台从 0 号编到 31 号。顾客手里的排队号其实就是内存地址——硬件把共享内存看成一条连续的长柜面，按 4 字节切成一个个“词”，再按词号分发到柜台上。

叫号规则只有一条：柜台号 = 词号 mod 32，也就是字节地址先除以 4 取整，再对 32 取余。为什么是 32？因为 32 正好对应一个 warp（线程束）里的线程数，硬件希望一个 warp 的 32 个线程在一个周期内各占一个柜台。映射公式如下：

把公式套到熟悉的矩阵上，问题就浮出水面了。假设我们放了一个 float m[32][32] 进共享内存，行主序存储下，m[row][col] 的字节地址是 (row*32 + col) * 4，词号是 row*32 + col，于是：

看明白了吗？同一列的所有元素，不管 row 变成几，bank 永远等于 col。也就是说，如果一个 warp 的 32 个线程恰好都去取第 0 列，32 个顾客会同时涌向同一个柜台。

小李恍然大悟：原来柜台分配是“地址说了算”，不是“线程说了算”。这就解释了为什么线程明明各自忙碌，却会在同一时刻挤进同一柜台。

bank=⌊addr4⌋ mod 32\displaystyle bank = \left\lfloor \frac{addr}{4} \right\rfloor \bmod 32
bank(m[row][col])=(row×32+col) mod 32=col\displaystyle bank(m[row][col]) = (row \times 32 + col) \bmod 32 = col
  1. 地址转词号把字节地址除以 4 并向下取整，把连续字节流切成 4 字节的单词，m[row][col] 的地址换算成“第几个词”。
    ⌊addr4⌋\displaystyle \left\lfloor \frac{addr}{4} \right\rfloor
  2. 词号落到柜台词号对 32 取模，得到 0~31 号 bank，柜台由地址决定而不是由线程决定。
    词号 mod 32\displaystyle \text{词号} \bmod 32
  3. 矩阵特例展开当行宽是 32 时，row*32 这部分的余数恒为 0，列号直接决定了 bank。
    (row×32+col) mod 32=col\displaystyle (row \times 32 + col) \bmod 32 = col

冲突类型

小李站在大厅里观察客流，看到了三种画风。第一种：32 个柜台各站一位顾客，一秒钟全部办完——这是理想状态，叫无冲突。第二种：几个柜台前各排了两人，需要开两轮业务才办完——这叫 2-way 冲突。第三种最夸张：31 个柜台空着，32 个人全堵在 0 号柜台前，要整整开 32 轮——这是最坏的 32-way 冲突。

硬件在每个时钟周期最多执行一次“内存事务”（一次事务 = 32 个柜台各服务一个词）。当多个线程命中同一个 bank 的不同地址时，它们必须排队，每轮只能进一位。所以一个 warp 的访存指令要分成多少轮，完全由“最挤的那个柜台”决定：

其中 $N_b$ 是命中第 $b$ 个 bank 的线程数。无冲突时每个 $N_b \le 1$，$T_{trans}=1$；全部挤进一个 bank 时 $T_{trans}=32$，访问延迟放大了整整 32 倍。这里有一个重要的例外：如果 32 个线程访问的是同一个地址（同一 bank 且同一词），硬件会启动“广播”——像前台复印一份公告同时发给所有人，一个周期就搞定，不产生冲突。冲突的判定标准不是“人数多”，而是“同 bank 不同地址”。

小李看着 0 号柜台前 32 人的长队，终于明白：再快的收发室，也架不住把所有人塞进同一个窗口。空着的 31 个柜台再闲，也帮不上忙。

Ttrans=max⁡0≤b<32Nb\displaystyle T_{trans} = \max_{0 \le b < 32} N_b
  1. 统计各柜台人数先数出每个 bank 被多少个线程命中，得到 $N_0, N_1, \dots, N_{31}$。
    Nb,b=0,1,…,31\displaystyle N_b,\quad b = 0, 1, \dots, 31
  2. 取最大值最挤的柜台决定了总共要开多少轮业务，一次共享内存指令就被拆成 $T_{trans}$ 次事务。
    Ttrans=max⁡bNb\displaystyle T_{trans} = \max_b N_b

诊断方法

怎么知道收发室堵没堵、堵得多严重？老工程师给了小李两样工具：一是站在大厅看监控（性能分析器），二是给每笔业务掐表（运行计时）。光靠“感觉好像有点慢”是没法向老板交代的。

专业做法是打开 Nsight Compute 之类的 profiler，它会在 Memory Workload Analysis 一栏直接给出共享内存负载（shared load/store）的 bank conflict 次数；旧一点的 nvprof 也提供 shared 存储效率指标。这些指标像大厅监控屏上的“各柜台排队人数”，一眼就能看出冲突出在哪个 kernel、哪条指令上。没有 profiler 的环境下，最朴素也最可靠的办法是A/B 对比：只改动 padding 一处，重编译、跑同一份数据，看内核总耗时是不是显著缩短。

做对比实验时，务必坚持“一次只改一个变量”，否则你分不清提速来自 padding 还是来自顺手删掉的那行断言。优化前后的耗时一除，加速比就出来了：

老工程师叮嘱：别猜，要量。拥堵不靠感觉，靠监控和秒表。小李拿着实测数据，底气十足地去改代码了。

S=TbeforeTafter\displaystyle S = \frac{T_{before}}{T_{after}}
  1. 优化前耗时记录未做任何改动时内核的整体运行时长，作为基线。
    Tbefore\displaystyle T_{before}
  2. 优化后耗时应用 padding 等优化后再次测量，得到新耗时。
    Tafter\displaystyle T_{after}
  3. 求比值两者相除得加速比，$S > 1$ 说明优化有效，$S$ 越大收益越明显。
    S=TbeforeTafter\displaystyle S = \frac{T_{before}}{T_{after}}

Padding优化

老工程师开出的第一张药方，是给每行末尾加一个“占位隔板”。想象大厅发号的方式：原来每行 32 个号码正好把 32 个柜台轮一遍，下一行又从 0 号柜台重新开始。如果大家都去取第 0 列，等于每行都拿 0 号，自然全挤进 0 号柜台。现在每行塞 33 个号码（多一个占位号），下一行的起点就整体错开了一位——第 0 行的第 0 列拿 0 号，第 1 行的第 0 列拿 1 号，第 2 行拿 2 号……人潮自动分流。

技术上的操作极其简单：把 float m[32][32] 改成 float m[32][33]，行尾多出的那一列就是 padding（填充元素，不存数据）。为什么加 1 就够？因为行宽 33 与 32 互质：33 mod 32 = 1，相邻两行在内存里整体错开一个 bank。重新计算 m[row][col] 的词号与 bank：

当 $W=32$ 时 $W+1=33$，上式变成 $(row + col) \bmod 32$。32 个线程同取第 0 列时，bank 分别为 0,1,2,…,31，各占一个柜台，一次事务全部办完。代价仅仅是多占 4% 的共享内存容量，换来的是最高 32 倍的访存加速，这笔买卖非常划算。

小李感叹：一块小小的隔板，成本几乎为零，却把大厅从“32 人挤一窗”变成了“32 窗各一人”。原来性能优化，有时就差这一列。

bank=(row×(W+1)+col) mod 32\displaystyle bank = \left( row \times (W + 1) + col \right) \bmod 32
  1. 行宽加一每一行末尾多一个占位元素，行宽从 $W$ 变成 $W+1$，物理内存多出一列但不存业务数据。
    W+1\displaystyle W + 1
  2. 重新计算词号行主序下，m[row][col] 的编号变为 row * (W+1) + col，padding 改变了跨行时的地址步长。
    row×(W+1)+col\displaystyle row \times (W+1) + col
  3. 取模求 bank对 32 取模。由于 $33 \equiv 1 \pmod{32}$，每往下走一行，起点正好错开一个 bank。
     mod 32,33≡1(mod32)\displaystyle \bmod 32,\quad 33 \equiv 1 \pmod{32}

消除策略

比“疏导拥堵”更高级的思路是“直接改变客流”——让顾客从一开始就不往同一个柜台走。老工程师给了三招，小李记在小本本上。

第一招：访问模式改写。 既然按列读会冲突，那就尽量按行读。如果业务上必须取某列的数据，先在共享内存里做一次“转置暂存”：用无冲突的行访问把数据写入一个临时 tile，再按行读出来用。这就像收发室把堆积的包裹先按楼层重新分拣一遍，让后续取件的人各自去不同柜台。

第二招：利用广播。 如果整个 warp 都需要同一个值，就让大家都读同一个地址。硬件对同地址读取有广播机制，一个周期同时发给 32 个线程，不产生任何冲突。好比大厅把同一张通知复印 32 份同时递出，比 32 个人分别查 32 份不同文件快得多。

第三招：权衡取舍。 padding 不是免费的：它增加共享内存占用，可能降低 block 的占用率（occupancy，即同时能驻留多少个 block）；而遇到 double 类型（8 字节）或多维数组时，+1 未必够用，要看访问步长。经验法则：行宽（步长）与 32 互质，列访问就不会冲突；步长是 32 的倍数，则必定冲突。所以分析时要先算步长与 32 的最大公约数。整体效率可以这样度量：

其中 $B_{used}$ 是这一回合实际被命中的不同 bank 数，$N_{threads}$ 是参与访问的线程数（通常为 32）。全部命中不同 bank 时 $E=1$，全部挤进一个 bank 时 $E=1/32$。

小李现在学会了先看监控、再算步长、最后决定加不加隔板、改不改客流。站在收发室门口，他已经是个合格的大堂经理了。

E=BusedNthreads\displaystyle E = \frac{B_{used}}{N_{threads}}
  1. 统计命中 bank 数一次内存事务中，至少有一个线程去访问的不同柜台数量。
    Bused\displaystyle B_{used}
  2. 除以线程总数用参与访问的线程数作分母，得到“柜台利用率”。
    Nthreads\displaystyle N_{threads}
  3. 效率解读$E=1$ 说明 32 个线程用满了 32 个 bank；$E=1/32$ 说明 32 个人挤在 1 个 bank 上，浪费了 31/32 的并行能力。
    E=BusedNthreads\displaystyle E = \frac{B_{used}}{N_{threads}}

AI 应用场景

矩阵转置 Kernel 优化

阶段：数据预处理 / 算子性能优化

解决的问题：图像与科学计算中频繁需要矩阵转置。朴素实现把数据分块搬进共享内存后按列写回，触发 32-way Bank Conflict，带宽被压到 1/32；用 Padding 让 tile 行宽 +1，列写回变成无冲突访问，吞吐数倍提升。

具体案例：某自动驾驶团队的 4096×4096 图像预处理转置算子，采用 __shared__ float tile[32][33] 后，内核耗时从 3.2ms 降到 0.6ms，整体预处理流水线延迟下降约 20%。

融合注意力 Kernel 中的 Softmax 归约

阶段：模型推理 / 融合算子

解决的问题：LLM 推理的 FlashAttention 类算子把 attention score 以 tile 形式暂存共享内存，计算 row-max 与 sum 归约时，跨行访问容易造成 2-way 到 4-way 冲突；通过 tile 行宽 padding 加一、并用广播读取共享标量，减少共享内存事务数。

具体案例：某 7B 参数模型推理引擎中，fused attention kernel 的共享内存 buffer 行宽 +1 后，kernel 耗时降低 18%，端到端首 token 延迟下降 7%，在不牺牲精度的前提下显著改善了用户体验。

Python 演示

python38 行
BANKS = 32  def bank_of(byte_addr):    word = byte_addr // 4    return word % BANKS  def simulate_access(shape, pad, kind, threads=32):    rows, cols = shape    width = cols + pad    hits = [0] * BANKS    for r in range(threads):        c = r if kind == "diag" else 0        word = r * width + c        b = word % BANKS        hits[b] += 1    return max(hits), sum(1 for h in hits if h > 0)  def report(name, shape, pad, kind="col"):    trans, used = simulate_access(shape, pad, kind)    print(f"{name:16s} 行宽={shape[1] + pad:2d} 事务={trans} 命中bank数={used:2d}")  def main():    print("== 32 线程同一时刻访问共享内存矩阵（模拟） ==")    print("\n-- 按列访问第 0 列 --")    report("32x32 无填充", (32, 32), 0)    report("32x33 Padding", (32, 32), 1)    report("32x34 Padding", (32, 32), 2)    print("\n-- 按对角线访问 --")    report("32x32 无填充", (32, 32), 0, "diag")    report("32x33 Padding", (32, 32), 1, "diag")  if __name__ == "__main__":    main()
目的：用纯 Python 模拟 GPU 共享内存的 Bank 分配规则，对比 padding 前后按列与按对角线访问的冲突度，直观验证“行宽与 32 互质则无冲突”的结论。
  1. word = byte_addr // 4每个 float 占 4 字节，把字节地址压缩成“词号”，对应硬件切词规则。
  2. return word % BANKS词号对 32 取模，模拟硬件把词分配到 32 个柜台的叫号规则。
  3. width = cols + padpadding 的本质就是让行宽多出若干列，从根本上改变跨行步长。
  4. c = r if kind == "diag" else 0用 kind 区分按列访问（都取第 0 列）与按对角线访问。
  5. word = r * width + c行主序存储下元素在内存中的词号，是判断 bank 的关键一步。
  6. hits[b] += 1给命中柜台计数，模拟一周期内“谁挤进哪个柜台”。
  7. return max(hits), sum(1 for h in hits if h > 0)事务数 = 最挤柜台的人数；第二项是实际用到的 bank 数。
  8. if __name__ == "__main__":保证脚本被 import 时不执行演示，直接运行时才跑。

常见误区

误区“我的程序访问全局内存很慢，是不是全局内存也产生了 bank conflict？全局内存和共享内存的优化不是一回事吗？”

正解Bank Conflict 是共享内存（以及局部内存等按 bank 组织的片上存储）特有的概念；全局内存的优化重点是合并访问（coalescing），它按 128 字节缓存行合并请求，而不是按 32 个 bank 并行分发。两者虽然都叫“内存优化”，但一个是“柜台分流”问题，一个是“整批装车”问题，不能混为一谈。

为何常见：两者都出现在访存性能优化的语境里，术语又都带“内存”二字，初学者很容易把不同层级的优化手段划等号。

误区“32 个线程同时读 `m[0][0]`，那就是 32 个人挤一个柜台，肯定 32-way 冲突，必须用 padding 绕开！”

正解硬件对“同地址”启用了广播机制：同一周期内多个线程读同一个词，数据被复制并发给所有线程，1 次事务就完成，完全不算冲突。冲突的严格定义是“同一 bank 内、不同地址”。就像同一条公告复印 32 份同时递出，比 32 个人分别翻 32 份不同文件快得多。

为何常见：从“柜台模型”做直觉外推时，很容易漏掉硬件专门为同地址访问设计的广播/多播电路。

误区“遇到冲突就加 1，padding +1 是万能药，所有矩阵问题都这么解。”

正解padding 的原理是让行宽与 32 互质，使跨行步长不被 32 整除。+1 只是最常见的一招（33 mod 32 = 1）；如果访问的是 double 数组或行宽本身有对齐要求，可能需要 +2 或重新设计布局。更微妙的是，padding 可能破坏原本无冲突的模式——比如对角线访问在 32×32 下无冲突，改成 32×33 后反而变成 2-way。正确的做法是先算步长与 32 的最大公约数。

为何常见：绝大多数教程只讲“32×32 加一列”这个经典案例，读者容易把特例当通法。

误区“共享内存延迟本来就比全局内存低很多，就算有冲突也就慢一点点，不值得为它费劲优化。”

正解32-way conflict 意味着一个本该 1 个周期的访存指令被拆成 32 个周期，在访存密集的内核里可能造成数倍乃至几十倍的整体减速。当然也要权衡另一面：padding 增加共享内存占用，可能降低占用率（occupancy），所以正确的姿势是“先测量、再优化、后回归”，而不是盲目堆 padding。

为何常见：共享内存确实比全局内存快，这种“快”的直觉容易让人低估冲突带来的相对损失——比例才是关键。

面试题

Q1. 现代主流 NVIDIA GPU 的共享内存被划分为多少个 bank？

单选题 · easy

  • A.8 个
  • B.16 个
  • C.32 个← 正确答案
  • D.64 个

答案：C。共享内存在硬件上被分成 32 个 bank，正好对应一个 warp 的 32 个线程。每个周期 32 个 bank 可以并行各服务一个 4 字节的词。就像收发室的 32 个柜台，数量是硬件设计定死的，不是 16 也不是 64。

Q2. 共享内存中每个 bank 每个时钟周期最多可服务的数据宽度是？

单选题 · easy

  • A.4 字节← 正确答案
  • B.8 字节
  • C.16 字节
  • D.32 字节

答案：A。每个 bank 的宽度是 4 字节，恰好容纳一个 float 或 int。32 个 bank 并行就是 128 字节/周期的理论带宽。8 字节是 double 占用的宽度，不代表 bank 服务宽度。

Q3. 对于 float m[32][32]（行主序），warp 内 32 个线程同时执行 s = m[threadIdx.x][0]，会发生什么？

单选题 · easy

  • A.无冲突，1 次事务
  • B.2-way 冲突，2 次事务
  • C.32-way 冲突，32 次事务← 正确答案
  • D.广播，1 次事务

答案：C。m[r][0] 的词号是 r×32+0，对 32 取模后全是 0，32 个线程全部命中 bank 0，属于不同地址、同一 bank，产生最严重的 32-way 冲突，需要 32 个周期才完成。这与“30 人挤同一个柜台”完全等价。

Q4. 若将上题的矩阵改为 float m[32][33]，线程仍读 m[threadIdx.x][0]，此时共享内存事务数是？

单选题 · mid

  • A.1 次← 正确答案
  • B.2 次
  • C.16 次
  • D.32 次

答案：A。词号变为 r×33，33 mod 32 = 1，所以 bank = r mod 32，32 个线程正好落在 32 个不同 bank 上，1 次事务完成。这就是 Padding 行宽加一的核心效果——每行起点错开一个柜台。

Q5. 为什么行宽从 32 加到 33 就能消除列访问冲突？

单选题 · mid

  • A.33 是质数更均匀
  • B.33 与 32 互质，跨行步长错开一位← 正确答案
  • C.行宽变大后容量更多
  • D.编译器会自动重排访问

答案：B。关键是互质关系：33 mod 32 = 1，相邻行在 bank 空间上整体偏移一格，列访问时 32 行依次落到 bank 0~31。容量增加只是副作用，编译器不会替你改访问模式，质数本身也不保证均匀。

Q6. 某 warp 内线程 t 访问共享内存词的索引为 16×t（步长 16 个词），冲突度为？

单选题 · hard

  • A.无冲突，1-way
  • B.2-way 冲突
  • C.4-way 冲突
  • D.16-way 冲突← 正确答案

答案：D。bank = (16t) mod 32，偶数线程落 bank 0，奇数线程落 bank 16，只有 2 个 bank 被使用，每个 bank 排 16 人，属于 16-way 冲突。判断方法：步长 16 与 32 的最大公约数是 16，能覆盖的 bank 数 = 32/16 = 2。

Q7. 关于共享内存“广播”特性的正确说法是？

单选题 · hard

  • A.同地址读取必须先串行排队
  • B.同地址读取会被硬件广播，不产生冲突← 正确答案
  • C.广播仅对写操作有效
  • D.广播会增加内存事务次数

答案：B。当多个线程访问同一个 bank 的同一个地址时，硬件把数据复制分发给所有请求线程，一次事务完成。这与“同一 bank 不同地址”的排队情形本质不同：前者是复印公告，后者是 30 个人分别查 30 份文件。

Q8. 下列哪些方式可以诊断共享内存 Bank Conflict？

多选题 · mid

  • A.Nsight Compute 的冲突计数← 正确答案
  • B.修改 padding 前后计时对比← 正确答案
  • C.查看共享内存的容量大小
  • D.手工阅读内核源码分析← 正确答案

答案：ABD。Nsight Compute 能在 Memory Workload Analysis 里直接给出冲突事务数；A/B 计时对比是朴素可靠的实验方法；读源码静态分析访存模式也能判断。查看容量只能知道用了多大空间，与冲突无关。

Q9. 下列哪些方法能有效消除或缓解 Bank Conflict？

多选题 · mid

  • A.矩阵行宽添加 padding← 正确答案
  • B.列访问改写为行访问或转置← 正确答案
  • C.增大 block 线程数量
  • D.利用广播让多线程读同地址← 正确答案

答案：ABD。padding 让跨行步长与 32 互质；改写访问模式从源头避开同 bank 拥挤；广播利用硬件特性把同地址读取变成一次事务。增大 block 线程数只是增加并行规模，访问模式不变，冲突依旧存在。

Q10. 关于 Bank Conflict 与 Padding 的更深入认识，正确的有？

多选题 · extreme

  • A.对角线下 32×33 反而 2-way 冲突← 正确答案
  • B.padding 一定减少共享内存占用
  • C.行宽与 32 互质可消除列冲突← 正确答案
  • D.padding 增占用但通常收益大← 正确答案

答案：ACD。对角线访问 m[r][r] 在 32×33 下词号为 34r，34 mod 32 = 2，只剩偶数 bank，变成 2-way，说明 padding 不是万能的；padding 会增加而非减少共享内存占用；行宽与 32 互质时列访问才能覆盖全部 bank；容量代价通常远小于冲突带来的几十倍减速，所以实践中往往收益更大。