3、无Conflict访问模式：连续访问与步长访问、Padding技术消除Bank Conflict、Swizzle模式介绍
好，咱们接着聊共享内存。上一节我讲了Bank Conflict是怎么一回事，说白了就是多个线程同时抢同一个Bank，导致访问串行化。那这一节，咱们就来看看怎么彻底避开它。
我个人习惯把无Conflict的访问模式分成三类来讲：连续访问、步长访问，还有Padding和Swizzle。这三种方式，基本覆盖了90%以上的实际场景。
3.1 连续访问：最理想的情况
先说说最简单的——连续访问。什么叫连续访问？就是线程0访问地址0，线程1访问地址1，线程2访问地址2……以此类推。
你想想看，共享内存有32个Bank，每个Bank的宽度是4字节（32位）。如果每个线程访问一个4字节的数据，而且地址是连续的，那正好每个线程落在不同的Bank上。完美！
核心规律： 地址 % 32 的结果决定了数据落在哪个Bank。连续访问时，这个余数刚好是0,1,2,...,31，一个都不重复。
我在项目中遇到过很多新手，上来就写这样的代码：
// 错误示范：每个线程访问同一个Bank
int idx = threadIdx.x;
float val = shared[idx * 32]; // 所有线程都访问Bank 0！嗯，这里要注意。上面这个写法，所有线程访问的地址都是32的倍数，余数全是0。32个线程全挤在Bank 0上，性能直接掉到1/32。
正确的写法应该是：
// 正确示范：连续访问
int idx = threadIdx.x;
float val = shared[idx]; // 每个线程访问不同的Bank说白了，连续访问就是最自然的访问方式。只要你的数据布局和线程索引对齐，Bank Conflict自然就不存在。
3.2 步长访问：小心踩坑
那如果步长不是1呢？比如步长是2、4、8？这就得小心了。
咱们先看步长为2的情况：
// 步长为2的访问
int idx = threadIdx.x * 2;
float val = shared[idx];线程0访问地址0（Bank 0），线程1访问地址2（Bank 2），线程2访问地址4（Bank 4）……看起来好像没问题？
等等，你算一下：32个Bank，步长为2，那线程0和线程16会怎么样？线程0访问Bank 0，线程16访问Bank 32，但Bank 32其实还是Bank 0（因为32 % 32 = 0）。
注意： 当步长是2的幂次时，很容易出现2-way、4-way甚至更高的Bank Conflict。步长2是2-way Conflict，步长4是4-way Conflict，步长8是8-way Conflict，步长16是16-way Conflict，步长32就是32-way Conflict（最坏情况）。
我曾经在一个图像处理项目里吃过这个亏。当时做矩阵转置，用了步长为32的访问模式，结果性能惨不忍睹。查了半天才发现是Bank Conflict搞的鬼。
那步长不是2的幂次呢？比如步长3、5、7？
这个就有意思了。因为3和32互质，步长为3时，线程0到线程31访问的Bank分别是0,3,6,...,93。你算一下93 % 32 = 29，刚好不重复。所以步长和32互质时，反而没有Conflict。
小技巧： 如果你必须用步长访问，尽量选和32互质的步长（比如3,5,7,9,11等）。这样能天然避免Bank Conflict。
3.3 Padding技术：简单粗暴有效
好，那如果步长就是2的幂次，改不了怎么办？比如做FFT或者某些特定算法，步长就是2、4、8。这时候Padding就派上用场了。
Padding的原理特别简单：在共享内存的每一行后面加一个或多个空位，把原本对齐的地址错开。
举个例子，假设你要声明一个32x32的共享内存数组：
// 不加Padding：每行32个float
__shared__ float shared[32][32];
// 加Padding：每行33个float，多出来的1个不用
__shared__ float shared_padded[32][33];不加Padding时，shared[0][0]和shared[1][0]的地址差32个float，也就是128字节。128 / 4 = 32个Bank，刚好对齐。所以shared[0][0]和shared[1][0]落在同一个Bank上。
加了Padding后，shared[0][0]和shared[1][0]的地址差33个float，也就是132字节。132 / 4 = 33个Bank，33 % 32 = 1。所以它们落在了不同的Bank上！
你看，就这么简单的一个技巧，把Bank Conflict给化解了。
经验之谈： Padding的大小怎么选？一般加1个就够了。但如果你遇到的是4-way Conflict，可能需要加更多。我个人的习惯是：先加1个试试，不行就加2个，最多加到4个。加太多会浪费共享内存，毕竟共享内存就那么点。
3.4 Swizzle模式：更高级的玩法
Padding虽然好用，但有个缺点：浪费空间。共享内存本来就金贵，每行多一个float，32行就多32个float，128字节呢。
那有没有不浪费空间的办法？有，Swizzle模式。
Swizzle说白了就是重新排列数据在共享内存中的存放位置。不是按自然顺序存，而是按某种映射关系存，使得访问时不会冲突。
举个例子，假设我们要做4x4的矩阵转置，每个线程处理一行：
// 原始数据布局（有Conflict）
// 行0: a00 a01 a02 a03
// 行1: a10 a11 a12 a13
// 行2: a20 a21 a22 a23
// 行3: a30 a31 a32 a33
// Swizzle后的布局（无Conflict）
// 行0: a00 a01 a02 a03
// 行1: a13 a10 a11 a12 // 循环左移1位
// 行2: a22 a23 a20 a21 // 循环左移2位
// 行3: a31 a32 a33 a30 // 循环左移3位为什么要这么移？因为转置时，线程0读第0列，线程1读第1列，线程2读第2列，线程3读第3列。如果不Swizzle，所有线程都读同一列，那就是4-way Conflict。Swizzle之后，每列的数据来自不同的原始行，就错开了。
当然，这个例子是简化版的。实际CUDA里的Swizzle更复杂，有各种模式，比如XOR模式、SHIFT模式等。
实战建议： Swizzle虽然高效，但实现起来比较麻烦。我一般只在性能要求极高、共享内存又特别紧张的时候才用。大多数情况下，Padding就够用了。
3.5 三种模式对比
最后，我整理了一个对比表格，方便你快速决策：
| 模式 | 优点 | 缺点 | 适用场景 |
|---|---|---|---|
| 连续访问 | 天然无Conflict，性能最好 | 要求数据布局和线程索引对齐 | 大多数通用计算 |
| 步长访问 | 灵活，适合特定算法 | 步长为2的幂次时容易出Conflict | FFT、卷积等 |
| Padding | 实现简单，效果明显 | 浪费少量共享内存 | 步长固定且无法改变时 |
| Swizzle | 不浪费空间，性能最优 | 实现复杂，调试困难 | 共享内存极度紧张的高性能场景 |
好了，这一节的内容就到这里。记住一句话：没有最好的模式，只有最合适的模式。下次你写共享内存代码时，先想想你的访问模式是什么样的，再决定用哪种技巧。
下一节我会讲更进阶的内容——如何利用共享内存做矩阵乘法优化。到时候咱们再细聊。