Shared memory bank 冲突：32 个 bank、串行重放与 padding 规避

更新时间：2026-09-25。前置阅读：GPU 内存层级、GPU 架构：warp 调度器。上一篇讲过，shared memory 是 block 内共享的片上 SRAM，延迟只有约 20~30 周期，是 GPU 分块优化的核心阵地。但这块高速存储有个隐蔽的脾气：它内部被切成 32 个 bank，一个 warp 的访问一旦在 bank 层面撞车，片上访问就会被拆成多轮串行执行，最坏慢 32 倍。本篇专门讲清 bank 冲突怎么发生、代价是多少、怎么绕开。

本文要回答的问题

• shared memory 的 32 个 bank 是怎么组织的，线程地址按什么规则落到哪个 bank？
• 无冲突、2 路冲突、32 路冲突和广播访问分别需要几个周期，延迟倍数是多少？
• 为什么矩阵转置和按列访问是 bank 冲突的重灾区？
• padding 一个元素为什么就能错开 bank，具体加在哪里？
• 除了 padding，交换索引顺序、调整分块排布等手法各自适用什么场景？

一、32 个 bank：每个周期各自只服务一个地址

shared memory 物理上是一块多体（multi-bank）SRAM，被均分为 32 个等宽的 bank，每个 bank 每个时钟周期可以独立响应一次访问，宽度通常是 4 字节（一个 32 位字）。32 个 bank 正好对齐一个 warp 的 32 个线程——这不是巧合，整个设计就是让一个 warp 的访问在理想情况下一拍做完。

地址到 bank 的映射规则非常简单：按 4 字节字编号顺序循环铺放，第 0 个字在 bank 0，第 1 个字在 bank 1，……，第 31 个字在 bank 31，第 32 个字又回到 bank 0。写成公式就是：

bank = (字地址) % 32        // 字地址 = 字节地址 / 4

当 lane k 访问 shared[k] 时，32 个线程恰好各进一个 bank，所有 bank 同时工作，一个周期（经流水后等效每周期一组）完成全部 32 次访问。这正是 shared memory 能每周期向寄存器馈送上百字节的原因。理解 bank 冲突的全部关键就是一句话：**限制不在总带宽，而在每个 bank 每周期只能打开一个地址。**线程访问不同 bank 彼此无关，访问同一 bank 的不同地址就得排队。

这种多体交叉的思路并非 GPU 独有，通用 SoC 里的内存交织也是同一原理，只是 bank 数和访问主体不同。

还要区分两个容易混在一起的指标：访问延迟和聚合带宽。无冲突时一次 warp 访问经流水发起，等效每周期可从 32 个 bank 各取一个 4 字节字，聚合峰值约 128 字节/周期/SM。这个「每周期」是稳态吞吐口径，不是第一次访问只要一个周期——shared memory 的取数流水本身有约 20~30 周期的延迟，流水填满后才逐周期吐出结果。bank 冲突伤害的也正是这条流水：本该一拍喂进去的一组访问被切成 n 拍，issue 端口被重放指令反复占用，不仅当前 warp 变慢，还挤占了同周期其他 warp 的发射机会。

bank 的宽度在历代架构上有微调。经典模型是 32 bank × 4 字节；Kepler 及以后允许在 4 字节与 8 字节 bank 宽度间配置，访问 double 或 long long 这类 64 位数据时切到 8 字节模式可避免一个数据横跨两拍；配置由程序在初始化时选定，映射公式相应改为按 8 字节字取模。bank 数量始终是 32，与 warp 宽度绑定，本篇其余数字均按最通用的 4 字节模型给出。

二、四种访问模式：并行、2 路、32 路与广播

一个 warp 对 shared memory 的任意访问模式，都可以按「每个 bank 被几个不同地址命中」分类。设某个 bank 被 n 个线程以 n 个不同地址命中，这 n 次访问无法并行，硬件把它们拆成 n 个串行事务（serial replay，线程依次重放取数），该 bank 上的延迟就是基准的 n 倍。整个 warp 这一轮访问的耗时，取决于冲突最严重的那个 bank。

访问模式	地址分布	串行事务数	相对延迟	有效吞吐
无冲突	32 线程落 32 个不同 bank	1	1 倍	100%（峰值）
2 路冲突	某 bank 被 2 个不同地址命中	2	2 倍	约 50%
n 路冲突	最重的 bank 被 n 个地址命中	n	n 倍	约 1/n
32 路冲突	32 个线程全落同一 bank 的不同地址	32	32 倍	约 3.1%（1/32）
广播读	32 线程读同一个字地址	1	1 倍	100%（数据广播给全 warp）

逐行看这张表。无冲突是基准，一拍完事。2 路冲突时，撞车的两个线程先服务一个、再重放服务另一个，这一轮整体耗时翻倍，其余 30 个 bank 被迫陪等，等效吞吐腰斩。n 路冲突按比例放大，比如一个 warp 内出现 4 路冲突，就按 4 个周期结算。最坏情况是 32 路冲突：所有线程访问同一 bank 内的 32 个不同地址，32 次访问逐一串行，延迟正好是无冲突时的 32 倍——原本约 20~30 周期的片上访问，退化到接近一次显存访问的量级，分块省下来的时间一次赔光。

广播是唯一的例外：当 32 个线程读的是同一个字地址时，硬件一次读出该字，通过广播网络把同一份数据发给全部 lane，仍然只算一个事务。程序里所有线程读取同一个 shared 常量（比如共享的半径、阈值、公共系数）时享受这条规则。注意广播只对读成立：多个线程写同一地址不是冲突，但写入结果未定义（除非各线程写的值恰好相同），不能指望广播语义。

把延迟折成绝对周期数更有体感。shared memory 基础访问约 20~30 周期，2 路冲突约 40~60 周期，4 路约 80~120 周期，32 路冲突约 640~960 周期——已经落在一次显存访问（约 400~800 周期）的区间。一个精心分块、本想把显存流量摊薄几十倍的 kernel，若读出阶段 32 路串行，等于每条数据又被拉回显存延迟的水平，分块收益清零。串行重放还会额外消耗指令发射槽和电力：每次重放都是一条真实的内存指令，占用 SM 的发射带宽，因此冲突严重时整个 SM 的所有 warp 都变慢，不只是肇事 warp 自己。

较新的架构（计算能力 7.0 及以后）还支持有限的多播（multicast）：同一 bank 内同一地址只需发给线程子集时，也可合并服务。但「不同地址必串行、同地址读可广播」的基本判据始终不变，本篇按最通用的 32 bank、4 字节模型讲解。

三、典型场景：按列访问与 stride 32

光看规则还差一层直觉，关键在 stride（步距）。设 lane k 访问的字地址为 base + k * stride，则它落入的 bank 是 (base + k*stride) % 32。

• stride = 1：bank 随 lane 均匀铺开，无冲突；
• stride 为奇数：与 32 互素，k 从 0 到 31 仍能遍历全部 bank，无冲突（如 stride 3、5 都安全）；
• stride 为偶数：只能覆盖部分 bank，出现多路冲突；
• stride = 32：k*32 % 32 恒为 0，32 个线程全部挤进 bank 0，标准的 32 路冲突。

一个 warp 内各 bank 的冲突程度可以不同：比如 lane 0、1、2 撞成 3 路，其余 lane 各自分散，整轮仍按最重的 3 拍结算——无冲突的线程并不能先走，它们的数据在同一事务里统一返回。因此优化时优先处理冲突路数最高的那小撮线程，收益最大；把一个混合模式从「3 路 + 大量无冲突」降到无冲突，整轮延迟直接从 3 拍回到 1 拍。非正方形矩阵的行宽不是 32 的倍数时，按列访问的冲突度还会随列号变化，有的列 32 路、有的列天然散开，这种「忽快忽慢」用取模公式逐列算一遍比凭感觉调参可靠得多。

stride 32 在 2D 数组里随处可见。C 语言中 float s[32][32] 的同一列元素 s[0][c]、s[1][c]、…、s[31][c] 行间距恰好是 32 个字。于是按列访问 = 32 路冲突，按行访问则无冲突。

矩阵转置把这个坑暴露得最彻底。朴素分块转置把 32×32 的源矩阵块搬进 shared memory 后，需要换一个方向读出：

__shared__ float tile[32][32];

// 合并加载：lane k 读 tile[r][k]，沿行展开，无冲突
tile[row][k] = src(...);
__syncthreads();

// 转置写出：lane k 沿列读 tile[k][col]，行间距 32
// 32 个线程全部命中 bank 0 → 32 路冲突，32 倍延迟
dst(...) = tile[k][col];

加载阶段沿行访问、风平浪静；转置读出阶段沿列访问，一拍能做完的事拖成 32 拍。这类代码最迷惑人的地方是逻辑完全正确、结果分毫不差，只有性能分析器里 shared memory 的吞吐曲线塌下去才看得见。按列处理的模板计算、以列主序排布的 Fortran 数据对接、直方图中对同组元素的聚集访问，都是同一类病根。

把完整的朴素转置 kernel 摊开看更清楚：

__global__ void transpose_naive(const float* in, float* out, int n) {
    __shared__ float tile[TILE][TILE];          // TILE = 32
    int x = blockIdx.x * TILE + threadIdx.x;    // 列坐标
    int y = blockIdx.y * TILE + threadIdx.y;    // 行坐标

    tile[threadIdx.y][threadIdx.x] = in[y * n + x];   // 沿行加载
    __syncthreads();

    int xT = blockIdx.y * TILE + threadIdx.x;
    int yT = blockIdx.x * TILE + threadIdx.y;
    out[xT * n + yT] = tile[threadIdx.x][threadIdx.y]; // 沿列读出：32 路冲突
}

最后一行是病灶：一个 warp 沿 x 维展开，threadIdx.x 从 0 跑到 31，访问 tile[0..31][c]，行间距 32，全部落 bank 0。注意冲突与全局坐标怎么算无关，只与线程在 shared 数组里跨没跨行有关。

归约（reduction）是另一个常见触发点。树形求和中半区线程读取 s[k + stride]，当 stride 是 2 的幂且数据按 32 列排布时，多个 lane 会以偶数步距同时命中同一批 bank，产生 2 路、4 路逐级变化的周期性冲突。这类冲突往往不致命（最多几路），但归约被调用数百万次时累计损失可观，逐行偏移或调整每段宽度即可消除。

还要区分清楚：bank 冲突是 shared memory 内部的事，和显存端的内存合并（coalescing）是两个独立机制。一次转置写出完全可能在显存端合并良好、却在 shared 读出时 32 路串行，两条腿要分别检查，参见 内存合并访问。

四、规避手法之一：padding 一个元素

最经典的解法出人意料地简单：把每行数组宽度从 32 改成 33。

__shared__ float tile[32][33];   // 每行多垫 1 个元素

行间距变成 33 个字后，lane k 沿列访问第 c 列时落入的 bank 是 (k*33 + c) % 32 = (k + c) % 32。k 从 0 到 31 依然遍历全部 32 个 bank，32 路冲突当场消失，代价仅是每 32 行浪费 32 个字（约 128 字节）的 shared memory——相对 164 KB/SM 的容量可以忽略，却换回 32 倍的延迟。

padding 的通用口径是：让行宽与 bank 数互素。33 与 32 互素所以有效；行宽若是 64、96 这类 32 的倍数，同样的灾难会原样重现。对更大的分块（如 64×16），垫一个元素通常也够用，动手前用上面的取模公式验证一遍最稳妥。

3D 数组和结构体数组套用同样的规则。float buf[D][32][32] 需要在每个 32 列平面上垫宽，即声明成 [D][32][33]；结构体数组则要算清楚一个结构跨几个字，padding 的位置永远是「让连续行（或连续 tile）的起点在 bank 上错开一个奇数」，而不是机械地在数组末尾加字节——末尾加字节不改变行间距，冲突原封不动。另外 padding 只影响 shared 内布局，显存里的源数组不必跟着改，两者在加载语句中天然解耦。padding 也不是没有副作用：每行多 4 字节，申请多个 tile 时浪费线性增长，极端情况下影响可驻留 block 数；另外行宽变成非 2 的幂，地址计算多一条加法指令，相对于消除的串行重放完全值得。

五、规避手法之二：交换索引顺序与重排映射

padding 改变的是存储排布，另一类手法改变的是「哪个线程访问哪个元素」的映射关系。

交换索引顺序是矩阵转置的自然写法：让加载时沿行展开的同一批线程，在写出阶段直接按转置后的坐标落位，即读 tile[k][c] 的地方改写成由线程索引互换后的配对线程去读 tile[c][k]。配合 __syncthreads() 保证写入对全块可见，shared 内始终只沿行访问，根本不进入列方向。这要求线程到元素的映射设计成双向都沿快变维展开，是分块转置 kernel 的标准结构。

**逐行偏移（对角线化）**用于无法 pad、或多个数组共享布局的场景：第 r 行的访问人为加上偏移 r，即访问第 c 列时实际取 (c + r) % 宽度。连续行在 bank 上逐行错开一位，和 padding 数学等价，只是把错开量从存储布局挪到了索引计算。归约、扫描这类沿列累加的 kernel 常用此法。

调整分块尺寸与数据类型打包则从源头避开 stride 32：把 tile 设计成 32×16、16×32 等形状使关键访问方向步距为 1；用 float4、double2 之类宽类型时，一次访问跨多个连续字，落在单个 bank 还是跨 bank 取决于架构的 bank 宽度配置（部分架构支持 8 字节 bank），换宽度后应重新核对映射。矩阵乘法分块中 tile 形状的选择，本质上就是同时兼顾显存合并、bank 冲突和 occupancy 的三方平衡。

索引交换落到转置 kernel 上，就是让写出阶段的读取仍沿行展开：块内线程以转置后的二维坐标（x、y 互换）去读 shared tile，实际读出的元素变成 tile[threadIdx.y][threadIdx.x]，配合写出时 x、y 坐标互换，shared 侧全程只出现行方向访问。与 padding 版本对比，它一行额外存储都不花，代价是线程映射设计稍绕，块尺寸不是正方形时坐标换算要格外仔细。工程上最常见的组合恰恰是两者一起用：tile[32][33] 加 padding 兜底，映射上再保证主访问路径沿行，双保险。

判断该用哪类手法有个简单顺序：先看能否重排线程映射避免沿列访问（零存储代价）；不能则加一行 padding（几乎零成本）；布局受接口约束再用索引偏移。无论哪种，改完都要用性能分析器的 bank conflict 计数复核——这个指标硬件直接给出，不必靠周期数猜。以 Nsight Compute 为例，关注 shared memory 板块的 bank conflict 相关计数（每 warp 平均冲突事务数）：无冲突 kernel 的该值应接近 1，32 路冲突时接近 32；数值是多少拍就对应几倍延迟，优化前后对比一目了然。

六、几个常见误区

• **以为 shared memory 快就自动零成本。**它的峰值延迟建立在无冲突前提上，32 路冲突直接把约 20~30 周期放大到约 640~960 周期，和显存延迟同档。
• **以为广播能救一切「相同 bank」。**广播只认相同地址的读；同一 bank 的不同地址照样串行，多线程写同一地址更是未定义行为。
• **只查显存合并不查 bank。**两者独立，转置类 kernel 常常合并满分、bank 冲突也满分。
• **padding 随手加、不验互素。**行宽必须与 32 互素才有效，64 宽 pad 到 65 有效，pad 到 48 无效。
• **忽视 occupancy 反噬。**tile 过大、padding 累积导致每块 shared 用量超限时，驻留 block 数下降，省掉的冲突延迟又从调度停顿里赔回去。
• 用 64 位数据还按 4 字节 bank 想当然。double 在 4 字节模式下横跨两个 bank、两拍取齐，换 8 字节 bank 配置或按 double2 打包后映射完全不同，数据类型一变就要重算。
• **只盯单个 warp。**bank 冲突按 warp 内 32 线程判定，不同 warp 之间互不干扰也不互相合并；块内不同 warp 的访问是否均匀，要分别检查，不能用整块的总访问模式下结论。

这一篇在知识体系里的位置

本篇是 GPU 硬件与编程模型系列的机制专题，上承 GPU 内存层级（shared memory 在金字塔中的位置与约 20~30 周期延迟）和 GPU 架构：warp 调度器（warp 32 线程的由来），与显存侧的 内存合并访问互为对照——一个管片上 bank 不撞车，一个管显存事务能合并。具体到应用手法，padding 与索引重排是 矩阵乘法分块中 tile 设计的必备检查项，矩阵转置则是本文规避手法的最小完整示例。

总结

shared memory bank 冲突的物理根源是 shared memory 被切成 32 个每周期只能服务一个地址的 bank，地址按「字编号 % 32」循环映射：warp 32 线程落不同 bank 时一拍并行完成（1 倍延迟、100% 吞吐），同一 bank 被 n 个不同地址命中就拆成 n 个串行事务按 n 倍延迟结算，2 路冲突即 2 倍延迟、吞吐减半，32 个线程全挤同一 bank 最坏慢 32 倍、吞吐只剩约 1/32，唯一例外是全线程读同一地址走广播、仍为一拍；stride 32 的按列访问和矩阵转置是典型重灾区，最有效的规避是把 tile 行宽从 32 垫到 33 做 padding（行宽与 32 互素即错开全部 bank），其次是交换线程索引顺序、逐行偏移和调整分块形状，改完务必用性能分析器的冲突计数复核，并与显存合并、occupancy 一起权衡。