访存优化：合并访问、bank conflict 与对齐

第 03 篇 说大多数算子卡在访存。但”卡在访存”还分两种：你能不能拿满该有的带宽，取决于访问模式。同样读同样多的数据，合并访问能跑满带宽，跨步访问可能浪费一大半。这一篇讲两个最重要的访问模式问题：global memory 的合并访问，和 shared memory 的 bank conflict。

一、合并访问：warp 的访存是按事务发起的

GPU 访问 global memory 不是一个线程一个字节地读，而是以 warp 为单位、按内存事务（transaction）发起。一个 warp 的 32 个线程同时发出访存请求，硬件把这些请求合并成尽可能少的、对齐的内存事务（通常以 32 字节或 128 字节为粒度）。

• 当一个 warp 的 32 个线程访问连续且对齐的 128 字节（每线程一个 float，正好 \(32 \times 4 = 128\) 字节）时，只需要少量事务就能取完，没有浪费。这叫合并访问（coalesced access）。
• 当线程访问分散（大跨步、随机），同样 32 个 float 的有用数据可能散落在很多条缓存行里，硬件被迫发起更多事务，每条事务里大部分字节是取来却没用的。有效带宽（你真正需要的字节 / 时间）随之崩塌。

用一个 gather kernel 实测：线程 t 读 in[t * stride]，stride 越大、访问越分散。RTX 3060 Ti 上 \(N=2^{26}\)、256 线程/block，统计有效带宽（只算真正需要的字节）：

stride=1 时拿到 412 GB/s，接近该卡 448 GB/s 理论上限；stride=32 时有效带宽只剩 90 GB/s，掉到约 1/4.6。原因不是 DRAM 变慢了，而是每条事务取回的 128 字节里只有 4 字节有用，其余被浪费——总线在搬运你不需要的数据。

实用规则：让相邻线程访问相邻地址。

// 好：相邻线程访问相邻地址，合并
int i = blockIdx.x * blockDim.x + threadIdx.x;
float v = in[i];

// 差：相邻线程跨大步，事务浪费严重
float v = in[i * 32];

处理二维数据时，这意味着让 threadIdx.x 对应行内连续维度（最内层、stride 为 1 的维度）。矩阵转置、某些 reduction 布局之所以慢，根因常常是访问跨了大步。解决办法通常是先用 shared memory 把数据”换布局”——但这又引出 shared memory 自己的访问陷阱。

二、bank conflict：shared memory 的 32 个并行通道

shared memory 为了高带宽，被分成 32 个 bank（存储体），每个 bank 每周期能服务一个 32 位字。32 这个数字和 warp 大小对齐：理想情况下一个 warp 的 32 个线程各访问一个不同的 bank，32 个请求一个周期全部完成。

bank conflict 发生在一个 warp 内多个线程访问同一个 bank 的不同地址时。硬件只能串行化这些访问：如果 32 个线程全落在同一个 bank（但不同地址），就要 32 个周期才能服务完，这叫 32 路冲突。

bank 的编号规则：地址（按 4 字节字计）模 32。所以 s[threadIdx.x]（stride 1）让 32 个线程落在 32 个不同 bank，无冲突；s[threadIdx.x * 32] 让所有线程落在同一个 bank，32 路冲突。

实测对照：让 block 内线程反复读 s[idx]，stride 控制冲突程度。这里只取无冲突（stride 1）和满冲突（stride 32）两端做对照，不展开中间 stride。RTX 3060 Ti，reps=100000：

32 路冲突慢约 11 倍。没有到理论 32 倍，是因为只有 shared 读这一步被串行化，循环里其他指令开销和延迟隐藏摊薄了差距；但一个数量级的代价足以说明问题。

padding 消除冲突

最常见的修复手法是加一列 padding。二维 shared 数组 s[N][32] 在按列访问时（s[i][threadIdx.x] 固定列变行）会产生冲突，因为同一列的元素地址间隔 32，全落同一 bank。把它声明成 s[N][33]，每行多一个无用元素，列访问的地址间隔变成 33（与 32 互质），32 个线程就散落到 32 个不同 bank：

__shared__ float tile[32][33];   // +1 padding 列，消除按列访问的 bank conflict

代价是多用一点 shared memory。矩阵转置、tiled GEMM 里这是标准做法，GEMM 篇 会再见到。

三、对齐与向量化访存

除了合并和 bank，两个补充手段能进一步压榨带宽：

• 对齐：global 访问最好对齐到事务边界（如 128 字节）。cudaMalloc 返回的指针默认对齐到 256 字节，所以从数组起点的连续访问天然对齐；但如果你从一个有偏移的位置开始访问（如 in + 1），可能错位、多取一条缓存行。

• 向量化访存：用 float4（一次读 16 字节）代替 4 次 float 读，能减少访存指令数、提高每事务的有效载荷。前提是地址按 16 字节对齐、数量是 4 的倍数。

// 向量化：每线程一次搬 4 个 float
const float4* in4 = reinterpret_cast<const float4*>(in);
float4 v = in4[i];
out4[i] = v;

向量化对访存密集 kernel（拷贝、element-wise、量化打包）常有可观收益，但要小心对齐和尾部不整除的处理。

四、把两件事串起来：转置为什么需要 shared

矩阵转置是访存模式问题的经典缩影：直接转置时，读是连续的（合并），写却是跨大步的（不合并），有效带宽被写端拖垮。标准解法是分块：

1. 一个 block 合并地把一个子块从 global 读进 shared；
2. __syncthreads()；
3. 从 shared 里按转置后的顺序读、再合并地写回 global。

shared 充当”换布局的中转站”，把两端的 global 访问都变成合并的。而 shared 里的转置读如果不加 padding，又会触发 bank conflict——于是 tile[32][33] 的 padding 正好补上。这个例子把本篇两个主题接到了一起，也是 kernel fusion 篇 里布局变换的基础。

五、小结与下一步

• global 访问以 warp 为单位按事务发起；相邻线程访问相邻地址才能合并，跨步访问让有效带宽实测从 412 跌到 90 GB/s。
• shared memory 分 32 个 bank，warp 内多线程访问同一 bank 不同地址会串行化；32 路冲突实测慢约 11 倍，padding（如 [N][33]）是标准修复。
• 对齐和 float4 向量化能进一步提升带宽利用。
• 转置类操作用 shared 当中转把两端访问都变合并，是这些技巧的综合应用。

访问模式之外，另一个影响带宽利用的因素是有没有足够多的并发 warp 来隐藏延迟。下一篇讲 occupancy 与延迟隐藏，以及为什么 occupancy 高不一定快。