共享内存能加速矩阵乘法，因为其延迟仅为全局内存的约1/10、带宽更高，通过分块（Tiling）使线程块协作预加载A行与B列，消除重复访存冗余；合理设置BLOCK_SIZE（如32）、规避bank conflict、必要时启用双缓冲，并用nsight-compute验证访存效率。
为什么共享内存能加速矩阵乘法
因为全局内存（Global Memory）带宽有限、延迟高，而共享内存是片上高速缓存，延迟只有全局内存的 1/10 左右，且带宽接近理论峰值。矩阵乘法中每个元素要反复读取 A 的行和 B 的列，如果不做优化，同一行 A[row][*] 会被同 block 内多个线程重复从全局内存加载几十次——这就是典型的**访存冗余**。共享内存的作用，就是让一个 block 内所有线程“合力搬运一次”，再本地共享使用。
共享内存分块（Tiling）的关键参数怎么设
核心是让每个 block 负责计算一个 BLOCK_SIZE × BLOCK_SIZE 的输出子块，同时把对应的一块 A 和一块 B 预加载进共享内存。常见陷阱是 BLOCK_SIZE 选太大导致寄存器溢出或 occupancy 下降：
-
BLOCK_SIZE = 16：适合小显存或老架构（如 compute capability 3.5），寄存器压力低，但 tile 太小，同步开销占比高 -
BLOCK_SIZE = 32：在 RTX 40/50 系列上最常用，平衡寄存器用量、shared memory 占用（32×32×sizeof(float) = 4KB）与 warp 利用率 -
BLOCK_SIZE = 64：需要至少 64KB shared memory / SM，现代 Ampere+ 架构支持，但容易触发 bank conflict，必须配合 swizzle 或 padding
别硬套数值——用 cudaDeviceGetAttribute(&val, cudaDevAttrSharedMemoryPerBlockOptin, dev) 查你卡的实际上限，再留 10% 余量。
共享内存访问必须避免 bank conflict
共享内存物理上分为 32 个 bank（Ampere 及以后仍是 32-bank），若两个线程同时访问不同地址但落在同一 bank，就会串行化，性能腰斩。典型冲突场景：
立即学习“C++免费学习笔记（深入）”；
- 二维数组按行主序声明：
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE]，线程(y, x)读As[y][x]→ y 相同、x 相差 32 的倍数时撞 bank - 没对齐的 float4 向量加载，也可能跨 bank
解决方法不是“不用二维数组”，而是改下标映射：As[threadIdx.y][threadIdx.x + threadIdx.y % 2]（简单奇偶偏移），或更稳妥地用 __shfl_sync + 寄存器暂存替代部分共享内存读取。
双缓冲（Double Buffering）什么时候必须上
当 K 维度很大（比如 K=8192），单次 tile 加载后计算时间远小于下一次加载准备时间，此时不重叠数据搬运和计算，GPU 就会空转。双缓冲本质是用两套共享内存 buffer 轮流工作：
- 第一轮：buffer_A 加载第 0 块 A，buffer_B 加载第 0 块 B → 计算
- 第二轮：buffer_A 加载第 1 块 A（同时 buffer_B 正在计算）→ 计算时 buffer_B 已就绪
关键点：__syncthreads() 不能放在整个循环末尾，而要拆成“加载后同步”和“计算后同步”；否则双缓冲退化为单缓冲。实测在 K > 2048 时，双缓冲常带来 15%～30% 提速，但代码复杂度明显上升，调试时注意 __syncthreads() 位置错位会导致随机结果错误。
真正难的不是写出来，是验证每一步访存是否真的被合并、bank conflict 是否清零、occupancy 是否拉满——这些得靠 nsight-compute 看 gld_efficiency、shared__inst_executed 和 sm__sass_average_data_bytes_per_sector_mem_shared_op_ld 这几个指标，而不是只看跑分。
