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

Shared Memory 是 CUDA 优化中性价比最高的手段之一。它位于 SM 片上，访问延迟比全局内存低一到两个数量级。

一、基本特性
1. 每个 block 独享，block 内线程共享；
2. 容量与 L1 共用一块片上 SRAM，每个 SM 可配置数十到两百多 KB；
3. 静态分配默认上限 48KB/block，超过需要用动态分配并调用 cudaFuncSetAttribute 设置；
4. 由 32 个 bank 组成，每个 bank 宽 4 字节。

二、Bank Conflict
同一 warp 内多个线程访问同一 bank 的不同地址时，访问会被串行化。典型例子是矩阵转置时按列访问，步长为 32 时全部冲突。解决方法是填充一列：__shared__ float tile[32][33]。

三、典型优化流程（以矩阵乘法为例）
1. 先保证全局内存合并访问，这是第一优先级；
2. 采用 tiling：每个 block 把 A、B 的子块加载到 shared memory，多次复用，可减少全局访存 TILE 倍；
3. 加载后调用 __syncthreads() 同步，计算完成后再次同步，避免覆盖；
4. 进一步使用寄存器分块（每线程计算多个输出）、双缓冲预取。
实测中，从朴素实现到 shared memory tiling 常见 3~5 倍提升。

四、常见坑
1. __syncthreads() 放在条件分支中，分支发散时会死锁或行为未定义；
2. 共享内存用量过大导致每 SM 驻留 block 减少，占用率下降；
3. 忽略 Tensor Core 对数据布局的要求，需要额外重排。

讨论：
1. 推理优化中，你们如何让 Tensor Core 与 shared memory 配合，用 CUTLASS 还是手写？
2. 分析 bank conflict 你们看 Nsight Compute 的哪些指标？

Shared Memory 还有个反直觉的点：bank conflict 只对 32 位访问才有意义，用 64 位或 128 位向量访问时，冲突计算方式完全不同，很多人照着 32 位的经验调，结果没效果。

补充 __syncthreads() 的一个坑：如果放在条件分支里，且分支在不同线程间发散，会直接死锁。这个坑我踩过，排查了半天。