4. 共享内存与Bank Conflict：共享内存的Bank映射、Conflict产生原因、Padding与Swizzle解决策略
共享内存，说白了就是GPU里的一块“高速缓存”，由程序员手动控制。它比全局内存快得多，延迟低、带宽高。但用好它，你得先搞明白它的“脾气”——Bank Conflict。
我个人习惯把共享内存比作一个多通道的仓库。仓库里有32个货架（Bank），每个货架同一时间只能服务一个工人（线程）。如果两个工人同时去同一个货架拿东西，那就得排队，这就是冲突。嗯，这里要注意，冲突不是错误，是性能杀手。
4.1 共享内存的Bank映射
共享内存被划分成32个Bank（对应Warp大小）。每个Bank的宽度是4字节（32位）。地址到Bank的映射规则很简单：Bank Index = (byte address / 4) % 32。
举个例子：
- 地址0-3 → Bank 0
- 地址4-7 → Bank 1
- 地址8-11 → Bank 2
- ……
- 地址124-127 → Bank 31
- 地址128-131 → 又回到Bank 0
你想想看，这就像循环队列。32个Bank循环映射地址。我在项目中遇到过，很多新手以为地址连续就能并行访问，结果发现性能没提升，就是因为没搞懂这个映射规则。
核心要点：同一个Warp内的32个线程，如果访问的地址映射到不同的Bank，就能并行访问，带宽翻32倍。如果映射到同一个Bank，就得串行化。
4.2 Conflict产生原因
Bank Conflict分三种情况：
- 无冲突（Ideal）：32个线程访问32个不同的Bank。完美！
- 部分冲突：多个线程访问同一个Bank，但访问的是同一个地址（广播模式，不算冲突）。或者访问不同地址，那就得排队。
- 完全冲突：所有线程都访问同一个Bank的不同地址。性能直接降到1/32。
为什么会这样？我举个例子：
// 假设共享内存 __shared__ float s[32][32];
// 线程tid访问 s[tid][tid] —— 这是对角线访问
// 每个线程访问不同的行和列，但地址映射到不同的Bank
// 无冲突
// 线程tid访问 s[tid][0] —— 所有线程访问第0列
// 地址：tid*32*4 + 0*4 = tid*128
// Bank Index = (tid*128/4) % 32 = (tid*32) % 32 = 0
// 所有线程都访问Bank 0！32路冲突！我曾经调试过一个矩阵转置的核函数，性能死活上不去。一查，就是这种列访问导致的Bank Conflict。32个线程抢同一个Bank，你说能不慢吗？
避坑指南：我曾经以为只有不同地址才算冲突，忽略了广播模式。实际上，如果多个线程访问同一个Bank的相同地址，硬件会广播，不算冲突。但如果是不同地址，就是真冲突。这个细节坑过我好几次。
4.3 Padding解决策略
Padding，说白了就是“加个垫片”。在共享内存的每一行末尾加几个空元素，让列访问的地址错开，避免映射到同一个Bank。
还是刚才的例子：
// 原始声明：__shared__ float s[32][32];
// 加Padding：__shared__ float s[32][32 + 1]; // 每行多一个float
// 现在访问 s[tid][0]
// 地址：tid*(33*4) + 0*4 = tid*132
// Bank Index = (tid*132/4) % 32 = (tid*33) % 32
// 因为33和32互质，所以tid从0到31，Bank Index会覆盖所有32个Bank
// 无冲突！我建议你记住这个技巧：Padding的大小通常选1或2个元素，保证行大小与Bank数互质。32的互质数有：1, 3, 5, 7, 9, 11, 13, 15, 17, 19, 21, 23, 25, 27, 29, 31。加1个元素变成33，就是最简单的选择。
个人经验：Padding不是万能的。如果访问模式特别复杂，Padding可能不够用。我一般先分析访问模式，再决定用Padding还是Swizzle。
4.4 Swizzle解决策略
Swizzle，说白了就是“重新排列”。通过改变数据在共享内存中的存储顺序，让原本冲突的访问变成无冲突。
Swizzle的核心思想：对地址进行位运算变换，重新映射Bank。常见的Swizzle模式有：
- XOR Swizzle：用异或运算重新排列行和列
- Shift Swizzle：用移位运算调整地址
- Custom Swizzle：根据具体访问模式自定义
举个例子，XOR Swizzle处理2D共享内存：
// 原始访问：s[row][col] → Bank = (row * width + col) % 32
// XOR Swizzle：s[row][col ^ row] → Bank = (row * width + (col ^ row)) % 32
// 假设width=32，访问s[tid][0]（列访问）
// 原始：Bank = (tid*32 + 0) % 32 = 0 → 冲突
// Swizzle后：Bank = (tid*32 + (0 ^ tid)) % 32 = (tid*32 + tid) % 32 = (tid*33) % 32
// 因为33和32互质，无冲突！你想想看，Swizzle比Padding更灵活。它不需要浪费额外的存储空间，但需要额外的计算开销。我一般在延迟敏感的场景用Padding，在带宽敏感的场景用Swizzle。
实战建议：Swizzle的位运算要小心。我遇到过因为异或操作写错，导致数据错乱的bug。调试了整整一天才发现是Swizzle逻辑反了。建议先在CPU上验证Swizzle的正确性，再移植到GPU。
4.5 知识体系图
下面这张图展示了共享内存Bank Conflict的核心知识结构：
4.6 实战对比
我整理了一个对比表，方便你选择策略：
| 策略 | 优点 | 缺点 | 适用场景 |
|---|---|---|---|
| Padding | 实现简单，无额外计算 | 浪费存储空间（约3%） | 列访问为主的场景 |
| XOR Swizzle | 不浪费空间，灵活 | 需要位运算，增加延迟 | 复杂访问模式 |
| Shift Swizzle | 适合特定步长访问 | 通用性差 | 固定步长模式 |
我的习惯：先分析Warp内线程的访问模式。如果是规律性的列访问，用Padding。如果是随机访问或者对角线访问，用Swizzle。记住，没有银弹，得根据实际情况选。
嗯，共享内存的Bank Conflict就讲到这里。说白了就是让32个线程别抢同一个Bank。Padding和Swizzle是两种主流解法，各有优劣。你在实际项目中多试试，慢慢就有感觉了。