Memory Bank 机制与加速计算

把数据从 global memory 搬进 shared memory，kernel 就一定更快吗？

刚开始学习 CUDA 时，这个想法很自然：shared memory 在芯片上，延迟低，还能让一个 block 的线程复用数据。但写出第一版 tile 后，常会碰到一个疑问：数据已经放进共享内存，为什么大家还是在等？

有时，问题不在数据离计算单元有多远，而在于同一条访存指令中，线程们是不是挤在同一个服务入口上。这个入口，就是 memory bank。

数据搬进来以后，怎样安排地址和线程，才能让共享内存真正发挥并行吞吐的优势？先从地址如何分配到 bank 开始。

目录

• Bank 为什么存在
• 地址怎样落到 bank
• 四种访问，四种结果
• 多一列，为什么能解开冲突
• Padding 之外的选择
• 哪些计算会遇到这个问题
• 把推导变成测量
• 最后记住的是地址
• 参考资料

Bank 为什么存在

假设有 32 个线程，各自要读一个数。如果存储只能一次服务一个请求，即使单次访问很快，也得排很长的队。把存储组织成多个可以并行服务的 bank，就有机会同时处理多个请求。

可以把 bank 想成图书馆的取书窗口，但要补上一条规则：窗口由书的地址决定，不由读者的身份决定。 同一个线程下一次换个地址，可能就去另一个 bank；多个线程也完全可能到同一个 bank。

Memory bank 是更广泛的存储组织概念，DRAM、片上 SRAM 等系统都可以使用 bank 来组织并行访问。不过，DRAM bank 有行激活、预充电等自己的约束，L2 缓存也有自己的组织方式。本文讨论的具体映射和冲突规则，聚焦于现代 NVIDIA GPU 的 shared memory，不能直接搬去解释 DRAM 或 L2。

在这里，两个「32」很容易被混在一起：

• 一个 warp 有 32 个 lane，lane 是线程在 warp 中的位置。
• Shared memory 有 32 个 bank，bank 是存储的组织方式。

它们分别属于执行端和存储端。不是 lane 0 永久拥有 bank 0，也不是给每个线程单独切一份共享内存。普通 shared memory 是 block 内线程协作使用的工作区，地址决定访问落在哪个 bank。NVIDIA Programming Guide：Shared Memory Access Patterns

Bank 的意义主要是吞吐：把可同时完成的访问分散出去。冲突会让请求拆成更多轮处理，但这不意味着所有 kernel 都会按冲突度等比例变慢。其他 warp、计算和流水线可能掩盖一部分等待。

地址怎样落到 bank

先只讨论最容易推导的情况：每个活跃 lane 对 shared memory 做一次标量、对齐的 32 位访问，例如读取一个 float。

常用映射是：

bank = (byte_address / 4) % 32

这里 / 是整数除法。连续的 4 字节字依次落到 bank 0、1、2……31，再绕回 bank 0。地址是交错分布的，不是先把一大段给 bank 0，再把下一大段给 bank 1。

若数组首地址对应 bank 0，float s[] 的元素 s[i] 就在 i % 32 号 bank。若首地址对应其他 bank，只会整体平移 bank 编号，不改变下面例子的冲突度。

图 1：每格是一个不同的 32 位字。图中省略 bank 3 到 30，省略处仍然按相同规律连续交错；数组首地址假设映射到 bank 0。本文配图均为根据地址公式绘制的示意图，不是硬件测量结果。

例如：

s[0]   → 字节偏移   0 → bank 0
s[1]   → 字节偏移   4 → bank 1
s[31]  → 字节偏移 124 → bank 31
s[32]  → 字节偏移 128 → bank 0

s[0] 和 s[32] 属于同一个 bank，但它们是不同地址。后面所有判断，都离不开这个区别。

四种访问，四种结果

假设一个 warp 的 32 个 lane 都活跃，lane 为 0 到 31，数据已经正确初始化。下面四个表达式各代表独立的一条标量读取，不是让同一个 kernel 同时执行四条后再混着分析：

// 地址分析示意。s 至少有 993 个 float；不是完整 kernel。
float a = s[lane];
float b = s[lane * 2];
float c = s[lane * 32];
float d = s[0];

第一个例子中，32 个 bank 都能参与服务。第二个例子中，lane 0 读 s[0]，lane 16 读 s[32]，它们都到 bank 0；其他偶数 bank 也发生同样的碰撞。第三个例子把 32 个不同字都压在一个 bank 上，丢掉了本来可以利用的并行服务能力。

第四个例子则很特别：大家要的是同一个数。硬件可以读取这个字，再广播给请求它的线程，不必为每个 lane 分别读取一次。NVIDIA Programming Guide：跨步访问与广播示例

图 2：每个地址标签都表示实际要读的字。冲突来自「同 bank、不同字」，广播来自「同一个字」。图只画代表性 bank，其余分布由文字说明。

对满 warp、不同地址、正整数步长的这类标量访问 s[base + lane * stride]，可以用一个小技巧检查推导：访问到的 bank 数是 32 / gcd(stride, 32)，冲突度是 gcd(stride, 32)。例如 stride 3 与 32 互质，虽然不是连续访问，仍会遍历全部 bank，没有冲突。

这是上述公式的数学推论，不是所有访存指令的通用检测器。部分 lane 活跃、广播、不同访问宽度或特殊指令，都要重新分析。

四个经常混淆的问题

Bank conflict 是服务能力的冲突，数据竞争是程序正确性的问题。 线程读取同一 bank 的不同地址，可以得到正确结果，只是处理需要更多轮。多个线程无同步地读写同一地址，则可能算错。所有 lane 读取 s[0] 的广播例外，不能推广成「所有 lane 同时写 s[0] 也安全」。普通并发写不能当作确定的合并或归约。

缓存 miss 关心数据在不在缓存里，bank conflict 关心请求怎样分配到 bank。 不能把 shared memory 的访问冲突称作「没有命中缓存」。

全局访存合并关心 warp 的 global 地址怎样组成内存事务。 Shared memory 的 bank conflict 发生在另一个层面。一次矩阵转置可以改善 global 读写的合并程度，却仍有严重的 shared memory 冲突。

同步负责建立执行顺序与内存可见性，不能改变 bank 映射。 __syncthreads() 能让 block 内消费者等到生产者写完；__syncwarp(mask) 适用于满足参与约束的 warp 内同步。它们不会把 stride 32 改成 stride 1。Warp convergence、CUDA stream 的任务排序，以及 MPI 的跨进程通信，也分别解决不同层级的问题。

为什么 double 和 float4 要另算

一个 double 占两个 32 位字，一个 float4 占四个。不能只把首地址代入公式，看到几个 lane 首地址碰到同一个 bank，就宣布整个 warp 有几路冲突。

一条宽访存指令可能天然需要多个事务或 wavefront；应按实际指令、每个事务覆盖的 lane 和字地址判断有没有额外冲突。数据量大而必需的多轮处理，与冲突造成的额外处理要分开。ldmatrix、异步拷贝、Tensor Core 相关布局还各有约束，源代码的一次访问也未必对应一条机器指令。Nsight Compute：Memory Workload Analysis

多一列，为什么能解开冲突

最经典的例子是一个 C/C++ 二维数组：

__shared__ float tile[32][32];

C/C++ 数组按行存储，最后一个下标变化最快。忽略首地址的整体偏移，tile[row][col] 对应的 float 下标是：

index = row * 32 + col
bank  = (row * 32 + col) % 32 = col

同一行的相邻列分散在不同 bank；同一列的 32 行却全部落在同一个 bank。

如果一个 warp 的 lane 0 到 31 各读 tile[lane][0]，下标分别是 0、32、64……992。这正是前面的 s[lane * 32]：32 个不同地址在 bank 0 排队。

现在只改数组声明：

__shared__ float tile[32][33]; // 每行尾部预留一个 float

逻辑上仍使用左侧的 32×32 数据，最后一列作为 padding。每行起点之间的距离从 32 个 float 变成 33 个：

index = row * 33 + col
bank  = (row * 33 + col) % 32 = (row + col) % 32

固定 col = 0，各行对应 bank 0、1、2……31；固定任意一行，连续 32 列也仍遍历全部 bank。Padding 没有增加 bank 数，它改变的是行与行之间的地址距离。

图 3：格内数字为 bank 编号，表格仅展示第 0、1、2、31 列与前 3 行、最后一行。右侧虚线格为第 32 列的 padding，不属于逻辑矩阵；其 bank 编号仍由真实地址决定。

放进矩阵转置里看

矩阵转置要把 input[row][col] 写成 output[col][row]。直接转置常让 global 读和写中的一侧出现大步长。Shared memory tile 可以在 block 内重排，使 global 两侧都采用连续访问；然后再处理 tile 内的列访问冲突。NVIDIA：An Efficient Matrix Transpose in CUDA C/C++

下面是一个完整的 kernel 定义，不含主机端分配与测试代码。输入为连续存储的 height × width FP32 矩阵，输出为 width × height 矩阵，两者不能重叠。启动 block 为 (32, 8)，grid 为 ((width+31)/32, (height+31)/32)：

__global__ void transpose_tiled(const float* input, float* output,
                                int width, int height) {
    __shared__ float tile[32][33];
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int x = blockIdx.x * 32 + tx;
    int y = blockIdx.y * 32 + ty;

    // 同一 warp 的 tx 连续：global 读与 shared 行写。
    for (int j = 0; j < 32; j += 8) {
        if (x < width && y + j < height)
            tile[ty + j][tx] = input[static_cast<size_t>(y + j) * width + x];
    }
    __syncthreads(); // 所有线程都经过这里，包括边缘 tile。

    x = blockIdx.y * 32 + tx;
    y = blockIdx.x * 32 + ty;
    // 交换 tile 坐标：shared 列读，global 连续写。
    for (int j = 0; j < 32; j += 8) {
        if (x < height && y + j < width)
            output[static_cast<size_t>(y + j) * height + x] = tile[tx][ty + j];
    }
}

blockDim.x = 32 时，同一 warp 的 ty 固定、tx 连续，所以这里的行写和列读符合前面的推导。CUDA 的线程线性编号以 x 维最快；若换成 (16, 16)，一个 warp 会跨两行，不能继续假定 threadIdx.x 等于 lane。

输出的边界条件同时保证对应输入元素已经加载，因此不会读取边缘 tile 中未初始化的格子。不要在同步之前让越界线程提前 return。

将 [32][33] 改成 [32][32]，可作为列访问冲突的对照版本。每个 block 的数组占用从 4,096 字节增到 4,224 字节，多了 128 字节。这只是容量变化；实际性能还取决于访存瓶颈和资源占用，需要通过基准测试比较。

Padding 之外的选择

Padding 很适合解释原理，但优化目标不是「凡是 32 都改成 33」。你可以从几种方向选择：

Swizzle 的思路是用确定的置换打散地址，而不是给矩阵添加空列。例如在一个 32×32 的标量 float 教学布局中，可把逻辑 (row, col) 存到物理 tile[row][col ^ row]。固定 row 时，异或是对 0 到 31 的排列；固定 col 时，bank 也遍历 0 到 31。写入和读取都必须使用同一个映射，否则会拿错数据。

这个例子仅由前面的 bank 公式推导，不能当作任意 Tensor Core kernel 的现成布局。实际 ldmatrix、向量加载与 TMA swizzle 还涉及粒度、对齐和特定指令的访问方式。支持 TMA 的架构有专门的硬件 swizzle 模式，需要遵守官方约束；相关功能能否使用，要以目标 GPU 的架构支持为准。NVIDIA：Asynchronous Data Copies 与 TMA Swizzle

Warp shuffle 则换了一个思路：若只是同一 warp 内交换寄存器里的值，可使用 __shfl_sync 系列，省去 shared memory 的写入、读取和相应同步。它不能直接交换其他 warp 的寄存器，也不会自动解决跨 warp 的数据依赖。对部分活跃线程，mask 和源 lane 必须满足指令的参与要求。NVIDIA：Warp Shuffle Functions

还有一笔资源账要算。Padding 或更大的 tile 可能使每个 SM 能同时驻留的 block 数减少，降低 occupancy；寄存器也可能成为限制。少一点冲突和少一点并发之间有取舍。某些 Tensor Core 布局原本依赖特定对齐，随手多加一列可能破坏向量化或合法性。最终要比较真实耗时。

哪些计算会遇到这个问题

矩阵转置最直观：global 两端希望合并访问，shared 中间要交换行列。上面的例子把这两层矛盾分别处理，适合作为第一个实验。

GEMM 的 tile更需要看具体线程映射。假设一条标量指令让所有 lane 读取 A_tile[row][k] 中同一个字，那可以广播；若各 lane 读取不同 row、固定 k，而行步长恰好是 32 个 float，则会冲突。不能仅看到「GEMM 使用 shared memory」就判断冲突，必须知道当前 warp 的每个 lane 正在取哪个元素。Tensor Core 实现还要继续追到实际加载指令。

规约和扫描容易在中间步骤出现跨步索引。例如一轮中活跃线程分别读取 s[2*lane]，就会重演 stride 2 的碰撞；随着活跃线程数和访问步长改变，每一轮的情况也会变。Warp 内可以考虑 shuffle，跨 warp 再用 shared memory 汇总。改变归约顺序还可能改变浮点舍入，应核对正确性。

FFT 的蝶形计算会在不同 stage 配对不同距离的元素。若某个 stage 的同一条 shared 读取呈现二的幂步长，就值得代入地址公式检查。别把不同时间发出的两条访问混成一次冲突，也别忽略复数或向量数据的访问宽度。已有库能满足需求时，先使用并测量 cuFFT 等成熟实现。

这些问题也会出现在 LLM kernel 中：GEMM、attention 的分块、softmax 归约，以及为矩阵指令准备数据的 shared 布局，都可能涉及 bank。但关注层级不同：一般框架使用者更应先看模型、batch、精度和后端选择；开发自定义 CUDA kernel 的人，才需要对具体 shared 指令逐 lane 推导。了解原理有助于读 profiler，却不会让一个普通框架参数自动消除 bank conflict。

把推导变成测量

最值得优化的情况是：目标 kernel 确实重要，而且分析发现 shared memory 的额外处理正在限制它。 若程序没有 shared 访问，或瓶颈在 PCIe 传输、global 带宽、计算、CPU 提交或其他同步上，调整 bank 布局的收益就很有限。

两个工具分工不同：

• Nsight Systems（nsys）看整个程序时间线：CPU、CUDA API、kernel、拷贝和 stream 的关系。先确定时间花在哪，是否有空隙或不必要的等待。官方 User Guide
• Nsight Compute（ncu）深入目标 kernel：查看 shared load/store、wavefront、冲突、占用率和关联源码。它的详细采集可能重放 kernel，不能把采集时的整段程序耗时当作业务基线。官方 Profiling Guide

选择工具时，要确认目标 GPU 和软件版本的支持范围。旧教程中的 nvprof 不支持 compute capability 8.0 及更高设备，这类设备应使用 Nsight 工具。Profiler Known Issues

下面的命令是实验流程示例，transpose_bench 是你自己的基准程序，并非本网站提供的可执行文件。基准程序需要分配输入输出、校验结果、预热、重复运行并记录 CUDA events：

# 保留源码与机器指令的行号关联；选择目标 GPU 对应的编译参数。
nvcc -O3 -lineinfo transpose_bench.cu -o transpose_bench

# 先看整体时间线。
nsys profile --trace=cuda,nvtx -o transpose-timeline ./transpose_bench

# 再采集目标 kernel，跳过前 10 次匹配的启动，采集一次。
ncu --section MemoryWorkloadAnalysis --section SourceCounters \
    --kernel-name 'regex:transpose_tiled' \
    --launch-skip 10 --launch-count 1 \
    -o transpose-detail ./transpose_bench

可按这个顺序做一次有结论的实验：

1. 测量：固定 GPU、工具与驱动版本、矩阵尺寸、数据类型和编译选项，先预热，用 CUDA events 重复测 GPU 时间，记录中位数与波动。正确性覆盖非 32 倍数尺寸。
2. 定位访问：在 ncu 的 Memory Workload Analysis 与 Source 页面找到主要 shared 指令，区分 load 和 store，并查看额外 wavefront。Wavefront 是硬件一次可处理的工作包，不是另一个 warp。
3. 推导映射：写出活跃 lane、字节地址、bank 和不同字地址的数量。确认线程线性编号与访问宽度，必要时查看 SASS，判断源代码实际变成什么指令。
4. 改布局：一次只改 padding、线程映射或 swizzle 中的一项；同步与边界也要一起保证正确。保存旧版作为对照。
5. 复测：先验结果，再比较独立基准中的耗时、shared 额外 wavefront 和 occupancy，最后看端到端时间。冲突下降而耗时不变，也是有效结论：它可能没有处在关键瓶颈上。

具体 metric 和 section 的可用性会随 GPU、ncu 版本变化，可用 ncu --list-sections、ncu --query-metrics 查询。不要把「wavefront 数大于 1」直接当成冲突，宽指令本来就可能需要多轮。汇总计数也可能包含 shared 与 L1 客户端之间的争用，应结合源码级证据核对。Nsight Compute：Compute Triage Guide

最后记住的是地址

Shared memory 提供了一个可以协作、复用、重排数据的片上工作区。Bank 让这个工作区能够并行服务；布局不合适时，程序又把访问挤回了队列里。

遇到「已经用了 shared memory，还是慢」时，可以先做一个很具体的动作：拿出其中一条读取，写下 lane 0、1、2、16、31 的地址，算一遍 bank。很多看起来神秘的冲突，会变成简单的余数运算。

然后回到 profiler：这条访问到底占多少时间？改完是否真的更快？把地址推导和测量连起来，才是 memory bank 能帮助加速计算的地方。

参考资料

• CUDA Programming Guide：Writing SIMT Kernels — warp、shared bank、跨步访问、广播及二维数组布局。
• CUDA C++ Best Practices Guide — shared memory 复用、冲突、资源占用与基于测量的优化。
• An Efficient Matrix Transpose in CUDA C/C++ — 转置中的全局访存与 padding；原文历史硬件跑分不代表现代 GPU 性能。
• Asynchronous Data Copies — TMA 与 swizzle 的架构、粒度及对齐约束。
• Nsight Compute Profiling Guide 与 Compute Triage Guide — 请求、wavefront、源码关联和瓶颈分析。
• Nsight Systems User Guide — 系统时间线与 CUDA tracing。
• 旧版 Profiler Known Issues — nvprof 的设备支持边界。