共享内存银行冲突由同一warp内多线程访问不同地址却落入同一银行触发，导致访问串行化；常见于行主序数组+列索引模式，需按GPU架构（32或64银行）针对性padding对齐stride，避免跨bank错位失效。
共享内存银行冲突是怎么被触发的
CUDA 共享内存被划分为 32 个（Compute Capability ≤ 7.5）或 64 个（CC ≥ 8.0，如 A100）独立银行，每个银行宽度为 4 字节。当多个线程在同一个 warp 内同时访问**不同地址但落在同一银行**时，就会发生银行冲突——访问被串行化，吞吐量直接下降。最常见诱因是数组按行主序连续布局 + 线程按列方式索引，比如 shared_mem[ty][tx] 在 tx 变化时地址步长为 1，极易撞进同一银行。
用 __syncthreads() 前先确认 bank 配置与访问模式对齐
银行数量不是固定 32；它取决于 GPU 架构和编译目标。例如：sm_75 默认 32-bank，sm_80 支持 32 或 64-bank 模式（由 cudaFuncSetAttribute 或编译器决定）。不查清实际 bank 数就盲目 padding，可能白忙活甚至恶化性能。
- 用
cudaDeviceGetAttribute(&val, cudaDevAttrSharedMemPerBlockOptin, dev)查当前设备最大共享内存容量，再结合 CUDA 文档确认 bank 数 - 避免“统一 padding 4 字节”：在 32-bank 下，4 字节对齐已足够；但在 64-bank（A100）下，需确保 stride 是 64 × 4 = 256 字节的倍数才真正无冲突
- 用
nvprof --unified-memory-profiling on或 Nsight Compute 的 “Shared Memory” 折叠页观察l__shared__inst__cycles_opr__ops__avg和l__shared__inst__ops__avg比值，比值 >1 就说明存在 bank conflict
二维共享内存无冲突布局的三种可靠写法
核心原则：让同一 warp 中相邻线程访问的地址，在内存中跨过足够多银行。常用手段是调整第二维大小（padding），使行首地址天然错开。
-
32-bank 设备（如 GTX 1080、V100）：声明
__shared__ float tile[16][16 + 1]（+1 列 padding），线程(tx, ty)访问tile[ty][tx]。此时每行起始地址偏移为(16 + 1) × sizeof(float) = 68字节 → 68 mod 128（32×4）= 68，保证相邻行起始落在不同 bank 组 -
64-bank 设备（A100）：改用
__shared__ float tile[16][16 + 4]（+4 列），使 stride = (16 + 4) × 4 = 80 字节 → 80 mod 256 ≠ 0，仍不够；更稳妥的是tile[16][32]，stride = 128 字节 → 128 mod 256 = 128，满足错位要求 -
彻底规避二维索引风险：用一维映射 + 手动 offset，例如
int idx = ty * (TILE_SIZE + PAD) + tx，其中PAD按目标 bank 数算好；这样你能完全控制 stride，也方便做循环展开或向量化加载
别忽略 shared memory 生命周期和 bank 冲突的耦合问题
bank conflict 不只发生在数据读写阶段。如果一个 warp 先写入共享内存，紧接着调用 __syncthreads()，然后所有线程又集中读取同一 bank 的某个标志位（比如 done_flag），这也会触发严重冲突——哪怕数据本身布局完美。
立即学习“C++免费学习笔记（深入）”；
- 把控制变量（如
is_valid,count）单独放在共享内存开头或结尾，远离主数据区，避免和高频访问区域共 bank - 避免 warp 内广播式读取单个标量：改用
__shfl_sync(0xffffffff, val, 0)从 lane 0 广播，不走 shared mem - 在双缓冲场景中，两个 buffer 若都按相同 padding 布局，要确认它们的基地址差不是 bank 宽度的整数倍，否则切换 buffer 时会继承冲突模式
真正难的不是算出 padding 数字，而是把 bank 边界、warp 调度、内存 coalescing 和同步点全部放到同一帧里看——少盯一个环节，优化就打水漂。
