CUDA 开发中共享内存Shared Memory是优化访存的重要手段但它并不是一块简单的快速缓存。很多 kernel 明明已经把数据放进了共享内存性能却不升反降原因往往就是 bank conflict。Shared Memory Swizzling共享内存重排是一种通过调整数据在共享内存中的布局来消除 bank conflict 的优化手段。这篇文章从共享内存的物理结构讲起解释 bank conflict 的产生原理然后用一个矩阵转置程序对比普通布局、padding 和 XOR swizzle 三种写法最后介绍如何用 Nsight Compute 定位冲突以及设计 swizzle 布局时容易踩的坑。1. 为什么共享内存访问会变慢从 Bank 冲突说起1.1 共享内存的物理结构共享内存在 GPU 的每个 SMStreaming Multiprocessor内部容量通常只有几十 KB 到上百 KB但带宽远高于全局内存。它之所以快是因为硬件把共享内存划分成了 32 个 bank每个 bank 是独立的存储单元每个时钟周期可以服务一个 word 的访问请求。这个 word 的大小一般是 4 字节对应一个float或int。当 warp 里的 32 个线程同时访问共享内存时硬件事务逻辑会检查这 32 个地址落在哪些 bank 上。如果每个地址落在不同的 bank那么这一次访问可以在一个事务内完成。如果多个地址落在同一个 bank硬件就必须把这个请求拆成多个事务每个事务服务其中一个地址然后再合成结果。这个拆分的次数就是冲突路数。例如线程tid访问s[tid]地址从 0 到 31分别落在 bank 0 到 bank 31这是最理想的情况bank: 0 1 2 ... 31 地址: 0 1 2 ... 31但如果线程tid访问s[tid * 2]地址是 0, 2, 4, ..., 62取模 32 之后得到的是 0, 2, 4, ..., 30, 0, 2, 4, ..., 30。每个 bank 上有两个线程访问这就是 2-way bank conflict实际需要两个事务才能完成这次访问。如果 32 个线程访问同一个地址s[0]硬件会启动广播机制所有线程拿到同一个值这种情况不算冲突代价仍然是一个事务。所以判断是否冲突关键看是否有多于一个线程访问同一个 bank 的不同地址。1.2 为什么二维 tile 也逃不过冲突在实际 kernel 中共享内存很少被当成一维数组直接使用更多是二维 tile比如__shared__ float tile[32][32]。二维数组在内存里是行优先连续存储的地址计算方式是address row * row_stride colbank 编号就是address % 32。当行宽row_stride是 32 时bank (row * 32 col) % 32 col这意味着 bank 只由列索引决定与行索引无关。如果 warp 内线程访问的列索引相同那么无论行索引怎么变化所有线程都会落在同一个 bank 上。这正是很多二维 tile 程序性能低下的根源。看起来tile[row][col]很自然但对硬件的 bank 分布来说可能存在严重的访问集中。矩阵转置就是一个典型例子后面会用完整代码演示。1.3 冲突路数的代价估算共享内存访问总延迟大致等于无冲突延迟乘以冲突路数。一个 32-way conflict 会把一次访问变成 32 次事务共享内存的吞吐优势基本被抹平。访问模式bank 分布冲突情况实际事务数s[tid]0..31 各一个无 conflict1s[tid * 2]0,2,4,...,30 各两个2-way2s[tid / 32]bank 0 被 32 个线程访问32-way32s[0]广播无 conflict1s[(tid ^ 7)]0..31 的排列通常无 conflict1实际项目里不能只凭感觉判断需要用性能分析工具统计共享内存的冲突次数。优化前先确认冲突存在优化后再确认冲突被消除这样每一步都有依据。2. 理解 Swizzling用重排布局消除 Bank 冲突2.1 Swizzling 的基本思想Swizzling 这个词在图形学里通常指对通道重新排列比如把 RGBA 变成 BGRA。在 CUDA shared memory 场景里它指对数据在共享内存中的物理位置做一次重排使得原本会落在同一 bank 的访问被分散到不同 bank。关键不是改变数据内容而是改变数据存放的“列位置”。比如原本应该存到tile[row][col]的数据现在存到tile[row][swizzle(row, col)]。只要swizzle函数是一个双射即每个有效的(row, col)都映射到唯一的列索引并且不同的(row, col)不会映射到同一个位置数据就不会丢失。为什么重排能消除冲突因为冲突的本质是“多个线程的地址取模后落在同一 bank”。原本 bank 只由 col 决定加入 row 的参与后bank 变成 col 和 row 的某种组合。warp 内线程如果 col 相同但 row 不同bank 也会不同冲突就被打散。2.2 Padding 是最简单的重排变体在讲 XOR swizzle 之前先看更常见的 padding 方案。把共享内存声明从tile[32][32]改成tile[32][33]行宽多一个元素__shared__ float tile[32][33];此时地址计算变成address row * 33 col bank address % 32 (row col) % 32因为33 % 32 1行号每增加 1bank 编号也增加 1。这样即使 warp 内线程的 col 相同只要 row 不同bank 也会分散开。Padding 的实现成本极低只需要改一行声明。缺点是浪费了一个元素的存储空间并且对于某些复杂的访问模式简单的单列 padding 不够灵活。比如行宽已经是 33 的倍数时冲突又会出现。2.3 XOR Swizzle用异或打散 bankPadding 的本质是把行号的信息叠加到列偏移上。XOR swizzle 也做类似的事但用的是按位异或。一个常见的最小公式是#define SWIZZLE(col, row) ((col) ^ (row))存储时不再写tile[row][col]而是写tile[row][SWIZZLE(col, row)]。为什么 XOR 有效因为异或有一个特性对于一个固定的row值col ^ row在col取遍 0 到 31 时结果仍然是 0 到 31 的一个排列。也就是说warp 内如果 col 连续变化、row 固定那么实际访问的 bank 编号恰好是 0 到 31 的乱序排列每个 bank 只被访问一次没有冲突。需要注意col ^ row的结果必须落在有效列索引范围内。如果 tile 宽度是 32col是 0..31row最好也只取低 5 位参与异或否则结果可能超出列范围。在 32x32 的 tile 里行号和列号都在 0..31所以可以直接用完整坐标异或。更通用的写法是#define SWIZZLE(col, row) ((col) ^ ((row) 31))这样无论 row 多大结果都保持在 0..31。XOR swizzle 不增加额外存储适合对共享内存容量敏感的场景。它和 padding 并不是互斥的有些高性能库会同时使用 padding 和 XOR swizzle进一步打散特殊访问模式。3. 实战矩阵转置的三种写法与性能对比矩阵转置非常适合用来演示 bank conflict。它既要用到二维 tile又要让线程从共享内存的列方向读取数据访问模式会直接暴露冲突问题。下面三个 kernel 分别对应普通布局、padding 和 XOR swizzle。3.1 环境准备和编译命令本文代码基于 CUDA C假设已经安装好 CUDA Toolkit并有一块支持共享内存的 NVIDIA GPU。代码里使用 32x32 的 thread blockN 取 1024便于演示。编译命令nvcc -archsm_80 -O2 -o transpose transpose.cu-archsm_80对应 Ampere 架构。如果你的 GPU 是其他架构换成对应的 compute capability比如sm_75或sm_90。编译成功后程序会依次运行三个 kernel 并输出耗时。3.2 普通布局会产生 32 路冲突的基线版本#include cuda_runtime.h #include cstdio #include cstdlib const int TILE 32; __global__ void transpose_baseline(const float* in, float* out, int n) { int tx threadIdx.x; int ty threadIdx.y; __shared__ float tile[TILE][TILE]; int row_in blockIdx.y * TILE ty; int col_in blockIdx.x * TILE tx; // 写入阶段tile[ty][tx]warp 内 tx 连续变化bank 不冲突 tile[ty][tx] in[row_in * n col_in]; __syncthreads(); // 读取阶段tile[tx][ty]warp 内 tx 连续变化但列 ty 固定 // 地址 tx * 32 tybank ty所有线程访问同一个 bank int row_out blockIdx.x * TILE ty; int col_out blockIdx.y * TILE tx; out[row_out * n col_out] tile[tx][ty]; }主函数启动方式如下后面两个 kernel 使用相同的网格和线程配置int main() { const int N 1024; size_t bytes N * N * sizeof(float); float* h_in (float*)malloc(bytes); float* h_out (float*)malloc(bytes); for (int i 0; i N * N; i) { h_in[i] i % 7; } float* d_in; float* d_out; cudaMalloc(d_in, bytes); cudaMalloc(d_out, bytes); cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice); dim3 block(TILE, TILE); dim3 grid(N / TILE, N / TILE); cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); transpose_baselinegrid, block(d_in, d_out, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms 0; cudaEventElapsedTime(ms, start, stop); printf(baseline: %.3f ms\n, ms); cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost); cudaDeviceSynchronize(); cudaFree(d_in); cudaFree(d_out); free(h_in); free(h_out); return 0; }这个版本在读取tile[tx][ty]时会产生 32-way bank conflict。因为 warp 内ty相同tx从 0 变到 31地址是tx * 32 ty取模 32 后全部等于ty。3.3 Padding 版本一行声明消除冲突__global__ void transpose_padding(const float* in, float* out, int n) { int tx threadIdx.x; int ty threadIdx.y; // 行宽从 32 改成 33额外的一列用于打破 bank 对齐 __shared__ float tile[TILE][TILE 1]; int row_in blockIdx.y * TILE ty; int col_in blockIdx.x * TILE tx; tile[ty][tx] in[row_in * n col_in]; __syncthreads(); int row_out blockIdx.x * TILE ty; int col_out blockIdx.y * TILE tx; out[row_out * n col_out] tile[tx][ty]; }这里只改了一个地方__shared__ float tile[TILE][TILE 1]。读取tile[tx][ty]时地址变为tx * 33 tybank 变为(tx ty) % 32。warp 内tx从 0 变到 31ty固定所以 bank 依次变化不再集中到同一个 bank。3.4 XOR Swizzle 版本不浪费额外存储__global__ void transpose_swizzle(const float* in, float* out, int n) { int tx threadIdx.x; int ty threadIdx.y; __shared__ float tile[TILE][TILE]; int row_in blockIdx.y * TILE ty; int col_in blockIdx.x * TILE tx; // 存储时对列做重排 tile[ty][tx ^ ty] in[row_in * n col_in]; __syncthreads(); int row_out blockIdx.x * TILE ty; int col_out blockIdx.y * TILE tx; // 读取时使用对应的逆重排 out[row_out * n col_out] tile[tx][ty ^ tx]; }这个版本的核心是读写两侧的列索引都加入异或。存储时原位置(ty, tx)的数据放在列tx ^ ty读取时原位置(tx, ty)的数据要从列ty ^ tx取出。因为异或满足交换律所以tx ^ ty和ty ^ tx是同一个值写读一致。从 bank 分布看写入阶段地址是ty * 32 (tx ^ ty)bank 是tx ^ tywarp 内tx连续变化时是排列无冲突。读取阶段地址是tx * 32 (ty ^ tx)bank 是ty ^ tx同样无冲突。3.5 性能对比怎么看三个 kernel 都完成了同样的矩阵转置但共享内存访问行为不同。实际运行耗时不是固定的跟 GPU 架构、编译优化、N 的大小都有关系。下面表格描述的是相对行为而不是精确数值版本共享内存布局读共享内存时的 bank 行为额外存储普通tile[32][32]32-way conflict无Paddingtile[32][33]无 conflict每 tile 多 32 个 floatXOR Swizzletile[32][32]无 conflict无在自己的机器上测试时建议用cudaEvent记录耗时并多次运行取稳定值。普通版本会明显慢于另外两个版本padding 和 XOR 的差距通常很小。注意矩阵转置中写共享内存阶段的冲突和读阶段不一定相同。这里普通版本冲突发生在读阶段写阶段反而是无冲突的。分析时一定要把读写分开看。4. 用 Nsight Compute 定位 Bank 冲突并验证优化效果4.1 为什么需要性能分析工具单纯比较耗时只能看到优化有没有效果看不到冲突在哪。Nsight Compute 是 NVIDIA 提供的 kernel 性能分析工具可以统计共享内存 bank conflict 的次数、访存吞吐、指令数等指标。优化 shared memory 访问之前最好先用它确认冲突的来源和量级。4.2 运行 Nsight Compute编译可执行文件后用ncu命令分析某个 kernelncu --set full ./transpose--set full会采集大量指标输出耗时较长。如果只想看共享内存冲突可以用--metrics指定指标ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./transpose指标名称在不同 CUDA 版本下可能略有差异如果命令报错可以在 Nsight Compute 的图形界面里打开 kernel 的 Memory Workload Analysis 页面查看 Shared Memory 部分的 Bank Conflicts 一行。4.3 如何解读结果对于普通版本读共享内存的冲突计数会很高理论上每个线程的读取都产生 32-way conflict所以总次数约等于线程数乘以 32 的某种折算。padding 版本和 XOR swizzle 版本的冲突计数应该接近 0。如果两个优化版本的冲突计数不为 0优先检查以下几点检查点说明读写规则是否一致存储时用了tx ^ ty读取时必须用相同的规则取回是否越界swizzle 后的列索引必须小于 tile 列宽是否缺少__syncthreads()写 shared memory 后必须同步否则读到未写入的数据线程块维度是否符合预期需要保证 warp 内线程的某一维连续否则分析会复杂这里要注意ncu默认会重新运行 kernel所以不需要在代码里做特殊处理。如果程序依赖命令行参数可以放在ncu后边。4.4 正确性验证性能分析之前先确保转置结果是正确的。可以用简单的方式验证让每个元素等于row * N col转置后检查out[row * N col] col * N row。也可以在 CPU 上写一个朴素转置对比。两个优化版本都没有改变数据的逻辑位置只是在共享内存里的物理位置不同所以输出应该和普通版本完全一致。5. 设计 Swizzle 布局的通用方法和常见坑5.1 先分析访问模式再设计 Swizzle不是所有 shared memory 程序都需要 swizzle。在动手之前先把 warp 内 32 个线程访问共享内存的地址列出来计算每个地址的 bank 编号。如果 bank 分布已经均匀加 swizzle 只会增加额外计算。设计 swizzle 的步骤可以归纳为确定线程到坐标的映射方式通常是(threadIdx.x, threadIdx.y)。确定每次共享内存访问的表达式比如tile[A(tx,ty)][B(tx,ty)]。把地址表达式代入bank address % 32看冲突路数。如果某维固定导致 bank 集中在该维上叠加另一维的信息常用方法是 padding 或 XOR。用ncu验证优化前后冲突计数并对比耗时。对于 32x32 的 tilecol ^ row是一个安全有效的 swizzle。对于其他尺寸要保证异或的结果仍然落在列范围内。一个通用做法是只取坐标的低 5 位参与异或例如#define SWIZZLE_INDEX(col, row) ((col) ^ ((row) 31))这样结果始终在 0..31 内不会越界。但要注意如果 tile 列宽小于 32需要进一步使用掩码或者改用 padding。5.2 不同访存粒度下的 bank 计算前面例子访问的是float一个线程一次访问 4 字节正好对应一个 bank。如果访问double每个数据占 8 字节会跨越两个 bank。如果访问float4占 16 字节会跨越 4 个 bank。硬件会把这种访问拆成多个 phase冲突分析要按 4 字节 word 重新计算。一个常见的经验是对于向量化访问swizzle 公式需要作用在“向量索引”上而不是字节地址上。例如存储float4 tile[32][8]每行 8 个float4每个线程访问一个float4bank 计算要按float4的元素个数扩展。不能简单套用标量版本的col ^ row。5.3 常见的坑和排查路径坑 1只改读取不改写入如果存储时用tile[ty][tx]读取时却用tile[tx][ty ^ tx]数据位置对不上最终结果会错乱。swizzle 的核心是读写规则必须对称要么都用同一套公式要么保证逆映射正确。坑 2异或结果越界col ^ row的结果可能大于 31。例如col 0row 32结果就是 32。这种情况会访问到下一行去轻则数据错乱重则触发非法访问。建议在 helper 函数里统一处理并加注释说明 row 必须小于 32或者使用row 31。坑 3忽略了同步共享内存是 block 内共享的写入后必须__syncthreads()否则其他线程读到的可能是旧数据。因为 swizzle 改变了物理位置有些线程可能访问到其他线程刚写入的位置数据依赖更强漏同步的后果更隐蔽。坑 4对所有访问都套用同一个 swizzle同一个 kernel 里可能有多种访问模式比如一次读 tile 的一行另一次读 tile 的一列。它们对 swizzle 的响应不同。不要默认一个公式解决所有问题要分别分析。坑 5过早优化如果 kernel 的瓶颈不在 shared memory而在全局内存访问、计算量或同步开销swizzle 并不能带来收益。先用 profiling 工具看整体瓶颈再决定优化方向。5.4 可以复用的检查清单下面的清单适合在写或改 swizzle 逻辑时逐项确认是否列出了一次 warp 访问的完整地址集合是否计算了每个地址的 bank 编号swizzle 函数是否是双射会不会出现两个坐标映射到同一个位置swizzle 后的列索引是否始终小于 tile 列宽所有读取路径和写入路径是否使用一致的映射写共享内存后是否调用了__syncthreads()是否用ncu对比了优化前后的 bank conflict 计数是否验证了 kernel 输出结果与 CPU 参考实现一致是否在多个数据规模下测试过性能是否确认 shared memory 不是当前瓶颈5.5 扩展方向XOR swizzle 不止能用于矩阵转置。在卷积计算中stencil 回读共享内存时经常出现跨行访问在 FFT 中radix 合并需要按不同步长读取在reduce中部分归约也会遇到 bank 冲突。这些场景都可以用同样的思路把固定维度的 bank 集中问题通过异或或 padding 打散。更复杂的布局还有 8x8 或 16x16 的 tile swizzle会同时考虑 L2 缓存 line 和 shared memory bank 的对齐。这些在 CUTLASS 等高性能计算库中很常见。实际工程里如果用了 CUTLASS通常不需要手写 swizzle因为它已经封装好了。起步阶段建议先把共享内存的 bank 机制吃透再用 Nsight Compute 反复验证动手写的 swizzle。比记住一堆复杂公式更重要的是建立一套“看访问 pattern - 推导 bank 分布 - 设计重排 - 用工具确认”的排查方法。这也是 CUDA 优化里最值得沉淀的能力。
