3. GPU内存访问优化

在GPU上做矩阵运算，算力往往不是第一瓶颈，访存才是。一个线程块里几百个线程同时向全局内存取数，如果地址散乱，硬件会把一次请求拆成几十次独立事务，带宽利用率可能跌到十分之一。本节围绕四个相互咬合的知识点展开：合并访问决定全局内存事务数，共享内存提供片上可编程缓存，bank冲突决定共享内存能否满速工作，而矩阵转置是把这三者同时用上的经典案例。

一、合并访问：让一个warp的请求落进同一批缓存行

全局内存以32字节扇区、128字节缓存行为单位搬运数据。一个warp有32个线程，若它们访问的地址连续且对齐，硬件可把请求合并成少数几个128字节事务；若地址间隔很大，则每个线程各触发一次事务。判断标准很直接：同一warp内相邻线程的地址差应等于元素大小。

以行主序的 float A[M][N] 为例，按列读取 A[ty][tx] 时，相邻 tx 对应的是行号变化，地址跨度为 N*4 字节。当 N=1024 时，一个warp的32次访问分散在32条不同的缓存行上，产生32次事务，有效带宽仅约 1/32。改成按行读取，同样32个线程落在连续128字节内，只需1次事务。这就是“合并”与“非合并”的差距来源，也是后面所有优化的动机。

二、共享内存：把重复读取搬到片上

共享内存是每个线程块私有的片上存储，容量通常为48KB至164KB（依架构与配置而定），延迟约为全局内存的二十到三十分之一。它的价值在于复用：当块内多个线程需要同一批数据时，先协作加载一次到共享内存，再从片上反复读取。

典型写法是分块（tiling）。以 C=A×B 为例，块大小取 32×32，每个线程负责计算 C 的一个元素。若不使用共享内存，计算一个32×32的输出块需要读取 32×32×32×2 个元素；使用共享内存后，只需加载两个 32×32 的输入块，读取量降为原来的 1/32。加载阶段用合并访问完成，计算阶段从共享内存取数，两个阶段各自发挥所长。

__shared__ float As[32][32];
__shared__ float Bs[32][32];
int tx = threadIdx.x, ty = threadIdx.y;
As[ty][tx] = A[row*N + k + tx];   // 合并：tx 连续
Bs[ty][tx] = B[(k + ty)*N + col]; // 合并：tx 连续
__syncthreads();
// 之后在循环中累加 As[ty][k] * Bs[k][tx]

注意 __syncthreads() 不可省略：加载与计算之间存在数据依赖，缺少同步会导致部分线程读到未写入的共享内存。

三、bank冲突：共享内存的隐形减速带

共享内存被划分为32个bank，每个bank宽4字节，地址按 addr/4 % 32 映射。一个warp的32个线程若访问落在32个不同bank上，可一次完成；若多个线程访问同一bank的不同地址，则串行化，这就是bank冲突。若访问同一地址，硬件会广播，不算冲突。

冲突最常出现在二维数组的列访问上。声明 __shared__ float tile[32][32] 时，tile[ty][tx] 的bank索引为 (ty*32 + tx) % 32 = tx，按行访问无冲突；但按列访问 tile[k][ty] 时，同一warp内 ty 变化而 k 固定，bank索引恒为 ty，32个线程恰好各占一个bank——这其实也无冲突。真正的问题出在跨步不为1的情形，例如 tile[ty][tx*2] 会让bank索引只取偶数，形成2路冲突。

通用解法是填充（padding）：把声明改为 __shared__ float tile[32][33]，使行跨度变为33。此时列访问的bank索引变为 (k*33 + ty) % 32 = (k + ty) % 32，随 ty 连续变化，冲突被打破。代价仅是多占一行空间。

四、矩阵转置的优化实现

把上述三点串起来，就得到标准的共享内存转置内核。设输入为 N×N 的 float 矩阵，块大小 32×32，网格按 N/32 划分。核心代码如下：

__global__ void transpose(const float* in, float* out, int N) {
    __shared__ float tile[32][33];          // 填充一列
    int x = blockIdx.x * 32 + threadIdx.x;  // 列
    int y = blockIdx.y * 32 + threadIdx.y;  // 行
    tile[threadIdx.y][threadIdx.x] = in[y * N + x];  // 合并读
    __syncthreads();
    x = blockIdx.y * 32 + threadIdx.x;      // 交换块坐标
    y = blockIdx.x * 32 + threadIdx.y;
    out[y * N + x] = tile[threadIdx.x][threadIdx.y]; // 合并写
}

关键细节有三处。第一，读入时 threadIdx.x 对应连续列，满足合并；写出时交换了 blockIdx.x 与 blockIdx.y，使同一warp的线程写入 out 的连续地址，同样满足合并。第二，tile 的第二维声明为33，正是为了在 tile[threadIdx.x][threadIdx.y] 这种转置读取中打散bank。第三，__syncthreads() 位于读写之间，保证整块数据加载完毕再转置。

若矩阵边长不是32的整数倍，需要做边界判断：if (x < N && y < N) 才写入，否则共享内存中会残留无效值并被写出。此时填充策略依然有效，但要注意越界线程不参与写回。

五、如何判断优化是否生效

不要只看总耗时，应结合指标定位瓶颈。用 Nsight Compute 观察 global_load_efficiency 与 global_store_efficiency：若低于80%，说明事务未充分合并；观察 shared_bank_conflict 相关计数，若存在明显冲突，检查共享数组跨度是否为32的倍数。一个可操作的验证流程是：先跑直接转置记录基线，再加入共享内存但不填充，最后加入填充，三次对比应能看到写效率与冲突计数的阶梯式改善。若加入填充后无变化，多半是访问模式本身已无冲突，或编译器把共享数组优化进了寄存器。

总结本节的主线：合并访问解决“全局内存事务数”，共享内存解决“重复读取”，bank冲突解决“片上并行度”，而矩阵转置恰好把三者放在同一个内核里检验。掌握这套组合，后续的卷积、归约、扫描等算子优化都能沿用同样的分析框架。