很多人刚接触 CUDA Shared Memory Swizzling 时，会觉得这是一个“高手专属”的优化技巧：反正 shared memory 已经比 global memory 快很多了，为什么还要费劲去改索引？我一开始也这样想。直到有一次写一个 32×32 的 shared memory tile 缓存，核心里面的计算量并不高，但 kernel 始终跑不到接近带宽的水平。最后用 profiler 看了一眼，问题根本不是计算，而是 shared memory 的访问在同一个 bank 上反复排队。
这个话题真正理解之后会发现，它不是一个孤立的骚操作，而是一套关于“地址如何映射到 bank、如何度量冲突、如何在可逆变换中摊开访问分布”的方法。这篇文章会把 bank 模型、padding、XOR swizzle、验证方法还有适用边界都拆开讲一遍，最后给你一个可以用到实际项目里的判断框架。
1. 先别急着优化：shared memory 的 bank conflict 是怎么拖慢内核的
1.1 shared memory 不是一个“随意访问”的高速数组
很多初学者会把 shared memory 想象成一个普通的高速内存：线程随便读写，只要数据在里面，速度就很快。实际硬件不是这样。在常见 NVIDIA GPU 架构中，shared memory 会被划分成一组 bank，通常是 32 个。每个 bank 在一个时钟周期内可以服务一次访问，而一个 warp 的线程会同时发出 shared memory 访问请求。
这个模型可以类比成 32 条并行的快递通道。如果 32 个线程分别去 32 个不同的 bank，那么一个周期就能全部完成。如果 32 个线程恰好打到同一个 bank，通道就只剩一条，其余线程只能排队。
所以 shared memory 快，不等于所有访问方式都快。快的前提是访问分布足够分散，让硬件并发处理。
1.2 bank conflict 的实际代价
当 warp 中的多个线程访问同一个 bank 的不同地址时，硬件会把这次访问拆成多次 wavefront。每次 wavefront 处理一批不冲突的访问。比如一个 warp 的 32 个线程全部访问同一个 bank 的不同地址，那就可能被拆成 32 次；如果不是所有线程都冲突，则拆成 2 次、4 次、8 次等等。
注意一个特例：如果多个线程访问的是同一个地址，很多架构支持广播，这不算典型的 bank conflict。真正麻烦的是“地址不同但 bank 相同”。这个细节经常被新手忽略，导致他们看公式时以为所有相同 bank 都会冲突。
bank conflict 的代价不是写代码的人能直接看到的。它不会报错，不会崩溃，只会让 kernel 变慢。如果你没有用 profiler 去看，可能还会把性能问题归咎于“GPU 不够快”或“算法复杂度太高”。
1.3 先判断：慢在计算，还是慢在访存
Shared memory swizzling 只对“访存冲突导致变慢”的情况有效。如果 kernel 慢是因为 Math 指令太多、因为 global memory 访存不合并、因为占用率太低，那么改 shared memory 索引不会带来什么改善。
我比较建议的做法是，在动代码之前先回答三个问题：
- kernel 的瓶颈是不是 shared memory load/store？
- profiler 里有没有明显偏高的 bank conflict 指标？
- shared memory 访问模式是不是有明显的周期规律，比如按行或按列固定偏移？
如果答案都是“是”，那 swizzling 才值得投入。否则，先优化更前面、更粗的瓶颈。
2. 从地址公式看 swizzling：一次可逆的索引置换
2.1 shared memory 地址如何映射到 bank
在大多数常见架构中，shared memory 的地址会按 4 字节宽度映射到 bank。对于一个float数组，偏移量是offset，那么 bank 可以近似看成：
bank = offset % 32如果你有一块float s[32][32]，并且按行主序存储，那么s[row][col]的偏移量是：
offset = row * 32 + col bank = offset % 32 = col也就是说，一个 32×32 的 tile，天然会按照列号落到 bank 上。这时候如果 warp 里的线程恰好是“同一列、不同行”的访问模式，它们就会命中同一个 bank，产生冲突。
这里要特别注意：不同架构的 bank 映射不一定完全相同，部分新架构还会引入地址散列。所以上面的公式适合用来理解概念和设计实验，真正落到某个具体 GPU 之前，应该以你当前 CUDA 版本对应的架构手册和 profiler 数据为准。
2.2 padding 是最朴素的“布局扰动”
最简单的避免方式不是 swizzling，而是 padding：把二维数组的宽度从 32 改成 33。
__shared__ float s[TILE][TILE + 1];这样s[row][col]的偏移量变成了：
offset = row * 33 + col bank = (row * 33 + col) % 32 = (row + col) % 32由于33 % 32 = 1，每换一行，bank 分布会整体平移一位。原本“不同行、同列”全部命中的情况，就变成不同 bank 了。
padding 的好处是简单、直观、不容易写错。坏处是每个 row 会多浪费一点空间，如果 tile 数量很多，可能影响 shared memory 占用和 occupancy。
2.3 XOR swizzle 是在索引层做一次位扰动
当我们希望既不额外占用 shared memory，又能改变 bank 分布时，可以在索引上做一个可逆变换。最常见的做法是 XOR swizzle。
假设一个 tile 的宽高都是 32，row和col都在[0, 32)内，可以这样计算线性索引：
#define TILE 32 __device__ __forceinline__ int tile_swizzle_index(int row, int col) { int r = row & (TILE - 1); int c = col & (TILE - 1); return r * TILE + (c ^ r); }这段代码的逻辑是：先把列号与行号的低 5 位做异或，再把最终结果映射到一块float s[TILE * TILE]的线性 shared memory 中。
此时 bank 会变成：
bank = (c ^ r) % 32如果 warp 中不同的线程按 row 变化、col 固定，那么c ^ r会随着 row 的 0 到 31 变化而产生一个完整的 32 项排列，冲突就被摊开了。
要注意的是，r * TILE这部分只是负责把不同 row 放到不同的地址区间，真正影响 bank 的是后面c ^ r的低 5 位。这个变换必须是一对一的，否则两个不同的(row, col)会映射到同一个 shared memory 地址，导致数据覆盖。
2.4 为什么不能随便加随机偏移
有人可能会想：既然要避免冲突，那给索引加一个随机数不就好了？不行。
Shared memory 的索引映射必须满足两个条件：
- 单射：不同的逻辑位置不能落到同一个物理地址。
- 可逆：需要能够从逻辑坐标算出物理地址，但不能出现两个坐标共享一个位置。
随机偏移很难保证这一点，而且还会让代码不可维护。padding 和 XOR 是两种“结构稳定、可验证、可解释”的做法，所以才会被反复使用。
3. 一个最小案例：从 baseline 到 padding，再到 XOR swizzle
3.1 用一份可 A/B 对比的示意代码
下面是一个接近框架的示意，不是完整业务 kernel，但可以用来理解如何替换索引映射：
#define TILE 32 __device__ __forceinline__ int plain_index(int row, int col) { return row * TILE + col; } __device__ __forceinline__ int padded_index(int row, int col) { int row_stride = TILE + 1; return row * row_stride + col; } __device__ __forceinline__ int swizzle_index(int row, int col) { int r = row & (TILE - 1); int c = col & (TILE - 1); return r * TILE + (c ^ r); }实际使用时，你只需要把 shared memory 的声明和索引函数配套：
__shared__ float tile_plain[TILE * TILE]; // 或 __shared__ float tile_padded[TILE * (TILE + 1)]; // 或 __shared__ float tile_swizzle[TILE * TILE];然后在内核里统一走：
int linear = swizzle_index(threadIdx.y, threadIdx.x); tile_swizzle[linear] = input_value;接下来，把“写入 shared memory”和“从 shared memory 读出”两段分别做 profiling。不要只看结果对不对，还要看访问模式。
3.2 三种方案放在一起看
| 方案 | 32×32 tile 的 bank 近似公式 | 额外 shared memory | 代码复杂度 | 适合情况 |
|---|---|---|---|---|
| 不处理 | col | 无 | 最低 | 小规模验证，冲突不明显 |
| padding | (row + col) % 32 | 每个 row 多 1 个元素 | 低 | 大多数二维 tile，代码好维护 |
| XOR swizzle | (col ^ row) % 32 | 无 | 中 | tile 宽为 2 的幂，想省 shared memory |
从实际项目角度看，我建议先上 padding。只有在 padding 导致共享内存占用明显增加、occupancy 下降，或者你已经确定 XOR swizzle 不会把代码搞乱时，再换 XOR。
3.3 一个容易忽视的问题：shared memory 大小也会影响结果
Swizzle 把32×32的 tile 存到 1024 个 float 里，看起来比 padding 的32×33省了 32 个 float。但要注意，shared memory 本身就是稀缺资源。
如果 block 数量很多，每个 block 都省 32 个 float，可能确实能提高 occupancy。但如果只是一个很小的 kernel，shared memory 占用远没到上限，这点节省就不重要。相反，XOR 代码更复杂，出错的概率更高。
所以选择方案时，不能只看“省不省 shared memory”，还要看“这个 kernel 是不是真的受 occupancy 限制”。
4. 不要只看核心里那几行：验证和排查链路很重要
4.1 一个可复用的四步判断流程
面对 shared memory bank conflict，我习惯用下面的顺序排查：
- 看现象：kernel 慢是整体慢，还是突然变慢？是计算型 kernel，还是访存型 kernel？
- 看指标：用 profiler 打开 shared memory 相关指标，确认 bank conflict 占比高不高。
- 看索引：把 shared memory 的访问抽象成
offset，手算几个关键 warp 的 bank 分布，判断是否有固定碰撞。 - 看全局：shared memory 优化之后，是否引入了新的问题，比如 occupancy 下降、global memory 不再合并、
__syncthreads()数量增加。
这套流程核心是“先量化，再改”。如果连指标都没看，就盲目套 swizzle，很可能把代码改复杂，性能却没有变化。
4.2 常见 metrics 怎么读
在 Nsight Compute 里，通常可以看 shared memory 相关的 bank conflict metrics。不同 CUDA 版本指标名称可能略有差异，但常见的是查看l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum和..._op_st.sum这类值。
这些指标只表示“发生了多少冲突”，具体是多大的冲突，还要结合访问指令数量来看。如果你的 shared memory 访问很少，即使冲突指标不为零，也可能不值得优化。
4.3 输出不对时，优先检查这三件事
Swizzle 之后如果结果不对，我的排查顺序是：
- 检查索引是否可逆：两个不同的
(row, col)是否映射到了同一个 shared memory 地址？ - 检查边界条件：
row和col是否越过了TILE范围？如果 tile 小于 32，row & (TILE - 1)的结果可能不是你想象的那样。 - 检查同步：写入 shared memory 之后有没有
__syncthreads()？在同一个 block 内，如果没有同步就读取，结果会不稳定。
另外还要注意，动态 shared memory 和静态 shared memory 的对齐方式不同：如果有extern __shared__，要谨慎计算每个 tile 的偏移量。Swizzle 本身解决的是 bank 冲突，不能替你解决越界和同步问题。
4.4 既然用了 swizzle，就要顺手把代码写清楚
我见过很多人在优化后留下一行看不懂的地址计算：return (row * 33) + (col ^ row);，既不写注释，也不解释为什么是 33，为什么异或。后面维护的人完全不敢动。
更合理的做法是：
- 把索引函数单独抽出来；
- 在函数上方注释清楚“这是为了消除 row 固定、col 变化时的 bank conflict”；
- 保留一个 baseline 版本，用宏或模板参数切换；
- 在 README 或注释里记录当时使用的 CUDA 架构和 profiler 数据。
这样即使几个月后再看，也能很快理解当初为什么这样做。
5. 什么时候该用 swizzling，什么时候不该硬上
5.1 适合用 swizzling 的场景
从经验看，下面这些场景比较适合做 shared memory swizzling：
- 二维 tile 的宽恰好是 2 的幂，比如 32、64。此时 XOR 公式简单，性能收益明显。
- 同一个 shared memory tile 会被多次读取，且读取方向既有行又有列，容易形成规律性冲突。
- profiler 明确显示 bank conflict 是热点，并且你已经排除 global memory 和计算指令的问题。
- shared memory 占用很紧张，不希望 padding 额外浪费空间。
满足这些条件时，XOR swizzle 是一个性价比不错的方案。
5.2 不适合硬上的场景
反过来，下面这些情况我不会急着用：
- 只是学习或验证功能：先跑通，再优化。不要一边学 API，一边引入索引变换，容易分不清是功能问题还是优化问题。
- tile 宽度不是 2 的幂：XOR 的位运算会变复杂，padding 往往更合适。
- shared memory 访问不是瓶颈：先优化 global memory 合并访问、减少重复加载、提高数学指令并行度，效果通常会更好。
- 已经用了 warp shuffle：如果能用
__shfl_sync避免 shared memory 访问，那就不需要额外 swizzle。 - 项目里没有 profiling 条件：没有量化手段，swizzle 就是盲改。改完你很难判断是真的变快了，还是冲突转移了。
5.3 长期维护：封装、记录、回退
Shared memory swizzling 本质上是一种“用算法复杂度换性能”的优化。它不会改变 kernel 的输入输出语义，但会改变内部布局。长期维护时，建议把布局策略当作一个模块来管理：
- 用
enum或常量表示LAYOUT_PLAIN、LAYOUT_PADDED、LAYOUT_XOR； - 将所有 shared memory 访问都通过
layout_index(row, col, layout)这个函数完成； - 在测试用例里对同一份数据分别跑三种布局，比较输出是否一致；
- 把性能对比结果记录到项目说明里，方便后续换卡时重新评估。
这样做的好处是，未来换 GPU 架构、换 CUDA 版本时，你可以快速重新 profiling，而不需要把整个 kernel 推倒重来。
最后的判断：先量化，再 swizzle
回到最开始的问题。CUDA Shared Memory Swizzling 并不是一个“能让 shared memory 变快”的魔法，它只是让 warp 在访问 shared memory 时尽量避免同一 bank 拥挤。它真正解决的问题，是把一个可预测的冲突访问模式，通过索引变换摊开成更均匀的分布。
如果只是凭直觉觉得“这个 kernel 慢，所以要 swizzle”，多半不会得到理想结果。比较稳妥的路径是：先跑通一个正确版本，再让 profiler 告诉你瓶颈在哪里，然后从 padding 开始试，最后才考虑 XOR swizzle。每一步都做量化对比，判断收益是否值得额外复杂度。
在 GPU 优化里，很多“高级技巧”本质上都是对资源访问模式的重新编排。Shared memory swizzling 也不例外。理解了这一点，你再看那些看起来奇怪的索引公式，就不会觉得它们高深，而是在看一段被注释掉的真实优化经验。
