CUDA Shared Memory：GPU 编程绕不开的那块片上缓存

一、先说痛点：GPU 快，但内存慢

GPU 的算力很强，可数据住在显存（global memory）里，访问延迟是几百个时钟周期量级。一个 kernel 如果每个线程都直接去 global memory 取数，SM 里的计算单元大部分时间都在等数据——算力再强也白搭。想榨性能，第一步永远是：让数据离计算单元更近。

二、Shared Memory 是什么：SM 片上的"小黑板"

每个 SM 上都有一小块片上存储，叫 shared memory，同一个线程块（block）内的所有线程都能读写，延迟只有几十个时钟周期量级，比 global memory 快一到两个数量级。你可以把它理解为程序员可控的 L1 缓存：数据放什么、什么时候换，全由你说了算。代价是容量小（几十 KB 量级），得省着用。

三、经典用法：分块（Tiling）

以矩阵乘法为例：把大矩阵切成小块（tile），每个 block 先把自己要用的两块 tile 一次性搬进 shared memory，之后的上百次乘加全在片上完成。数据复用率越高，global memory 访问被摊薄得越厉害。这是 CUDA 优化里出镜率最高的一招。

四、最常见的坑：Bank Conflict

Shared memory 在硬件上被分成 32 个 bank，连续 4 字节分属不同 bank。同一拍里，如果 warp 中多个线程访问了同一个 bank 的不同地址，访问就会被串行化——这就是 bank conflict，性能悄悄掉一大截。注意：多个线程访问同一 bank 的"同一个地址"是广播，不算冲突，别误杀。

五、躲坑实战：补一列 Padding

最经典的场景是 tile 转置或按列访问时，每行 stride 恰好是 32 的倍数，导致一列数据全挤在同一个 bank 里。解法简单粗暴：声明数组时多补一列（比如 [32][33] 而不是 [32][32]），错开 bank 映射，冲突立刻消失。Nsight Compute 里能直接看到 bank conflict 计数器，优化前后对比一目了然。

六、结论

Shared memory 是 CUDA 编程从"能跑"到"跑得快"的分水岭：理解分块思想、躲开 bank conflict，很多 kernel 的性能提升是立竿见影的。记住一句话：数据复用一次，就赚一次 global memory 的延迟。

Shared Memory 性价比高是因为它把「带宽」和「延迟」一起改善了——它是片上存储器，访问延迟比全局内存低一两个数量级，而且带宽高。核心优化点：1）bank conflict 是性能杀手，32 个 bank 的访问模式要设计好，最典型的坑是按列访问二维数组（stride = 32 时全撞一个 bank）；2）__syncthreads() 要放在所有线程都会到达的位置，放条件分支里且分支发散会死锁；3）优化优先级上，全局访存的合并（coalesced）通常比 shared memory 更基础，先把全局访存理顺再优化 shared。实测矩阵乘法用上 shared memory 能有数倍提升，但继续堆 register blocking 收益就递减了，因为瓶颈转移到了别处。