8. 共享内存
基础概念
Bank conflict —— 以矩阵转置为例
其实之所以会有所谓的冲突与解决方法，和 硬件实现 是息息相关的。
硬件层面的 Bank Conflict 是如何发生的？
硬件上，这一大块共享内存，被切分成了 32 个独立的小 SRAM 模块，这就是 Bank。
- 独立工作： 这 32 个 Bank 拥有各自独立的地址解码器和数据总线。它们可以同时响应各自的读取/写入请求。
- 交叉编址： 内存地址被像发扑克牌一样，轮流分配给这 32 个 Bank（即地址 0 在 Bank 0，地址 1 在 Bank 1… 循环往复）。
在线程和这 32 个 Bank 之间，硬件中布置了一个极其复杂的电路网络，叫做 Crossbar Switch（交叉开关）。它的作用就像一个超级电话交换机，负责在瞬间将 32 个线程的请求路由到对应的 Bank 上。
当 Warp 中的 32 个线程发出内存访问请求后，Crossbar Switch 会瞬间解析这 32 个地址。
- 理想情况（无冲突）： 32 个地址刚好散布在 32 个不同的 Bank 中。Crossbar Switch 完美连线，32 个 Bank 同时发力，一个周期内 32 个数据全部取出！
- 冲突发生（Bank Conflict）： 假设线程 0 要读 Bank 0 的地址 A，线程 16 要读 Bank 0 的地址 B。Crossbar Switch 发现有两个请求都指向了 Bank 0。但是！Bank 0 在物理上仍然是一个单端口 SRAM，它在一个周期内只能解码一个地址。
- 硬件的无奈之举（串行化）： 面对指向同一个 Bank 的多个不同地址请求，Crossbar 的硬件仲裁器（Arbiter）只能强制介入，把这些请求打散成多个时钟周期来执行。有几个请求挤在同一个 Bank，就要耗费几个周期（这就是 n-way conflict）。
《Unifying Primary Cache, Scratch, and Register File Memories in a Throughput Processor》
在论文的第 3 节（Baseline GPU Architecture）
“Each SM contains 64KB of cache and 64KB of shared memory. Each of these structures is composed of 32 2KB banks, and each bank supports one 4-byte read and one 4-byte write per cycle… Shared memory supports scatter/gather reads and writes, subject to the limitation of one access per bank per cycle.”
解读：这段官方描述完美印证了我之前的解释：共享内存由 32 个独立的 Bank 组成，而每个 Bank 在一个时钟周期内只能支持一次读操作（one read per cycle）。如果有多个读请求发向同一个 Bank，必然无法在一个周期内完成，这就是 Bank Conflict 的物理根源。
为什么不直接造出完美的内存？
你可能会问：“既然单端口 SRAM 会导致排队，那直接造一个拥有 32 个读写端口的 SRAM 不就永远没有 Bank Conflict 了吗？”
答案很现实：成本、面积和功耗。
在集成电路设计中，给 SRAM 每增加一个端口，其内部连线复杂度和所占用的硅片面积几乎呈平方级增加。一个真正的 32 端口 SRAM，它的体积和发热量会极其恐怖。如果 GPU 采用这种设计，原本能塞下上万个流处理器（CUDA Cores）的芯片，可能连几百个都塞不下了。
二维线程块，Warp 是怎么切分的？
无论你在代码里把 Thread Block 定义成 1 维、2 维还是 3 维，GPU 硬件在底层统统都会把它们“拍扁”成一个 1 维的线性队列，然后每 32 个线程切一刀，划分为一个 Warp。
GPU 拍扁多维数组的规则是：x 维度变化最快，然后是 y 维度，最后是 z 维度。（这和 C/C++ 中二维数组在内存中的按行存储是一样的）。
对于二维 Block，线程的线性索引（Linear Thread ID）计算公式为：
线性 ID = threadIdx.y * blockDim.x + threadIdx.x
说些什么吧！