内存访问模式基础：合并访问、内存事务与带宽利用率

核心原理：从“请求”到“事务”的硬件视角

GPU 的高吞吐量建立在海量线程并行执行之上，但每个线程最终都要通过内存系统获取数据。理解内存访问模式，首先要建立“内存事务（Memory Transaction）”的概念。GPU 内存控制器以固定大小的粒度处理请求，这个粒度通常为 32 字节（一个扇区）或 128 字节（一个缓存行）。当线程束（Warp，32 个线程）中的线程访问内存时，硬件并不会为每个线程单独发起一次内存请求，而是将整个线程束的访问请求合并为尽可能少的内存事务。

所谓合并访问（Coalesced Access），是指线程束内 32 个线程访问的地址落在同一个或少数几个内存事务覆盖的连续区域内。理想情况下，32 个线程访问 32 个连续地址（例如 float 数组的 32 个连续元素），硬件只需发起 1 个 128 字节的内存事务即可满足全部需求。反之，若线程访问的地址分散在不同缓存行中，则可能触发多个事务，造成带宽浪费。

内存事务的粒度与硬件行为

现代 NVIDIA GPU（Volta 及以后架构）的 L2 缓存行大小为 128 字节，而内存子系统的最小事务粒度是 32 字节扇区。这意味着：即使线程束只访问一个 4 字节的 float，硬件也会拉取包含该地址的整个 32 字节扇区到 L2 缓存。如果线程束的 32 个线程恰好访问 32 个连续 float（128 字节），则正好覆盖 4 个扇区，构成一个完整的 128 字节事务，无任何浪费。

然而，如果线程束访问的是按 128 字节步长跳变的地址（例如 stride = 32 个 float），则每个线程都落在不同的 128 字节缓存行中，硬件需要发起 32 个独立事务，每个事务仅使用 4 字节有效数据，带宽利用率骤降至约 3%。这种模式通常被称为“跨步访问”或“散列访问”，是性能杀手之一。

方法流程：如何设计合并访问的索引映射

在实际编程中，合并访问的关键在于让线程索引与数据地址呈线性关系。以下是一个典型的矩阵转置案例，展示错误与正确的索引映射方式。

场景：将 1024×1024 的 float 矩阵转置。每个线程处理一个元素，线程束的 32 个线程应尽量访问连续的行元素。

错误写法（按列读取）：

// 每个线程读取 A[row][col]，但 col = threadIdx.x
// 线程束内 threadIdx.x = 0..31，访问 A[row][0..31]
// 但 A 按行存储，A[row][col] 的地址 = row * 1024 + col
// 线程束访问连续地址 → 合并 ✓
// 然而写入 B[col][row] 时，地址 = col * 1024 + row
// 线程束内 row 相同，col 不同 → 写入地址间隔 1024*4B = 4KB
// 写入非合并 ✗
__global__ void transpose_bad(const float* A, float* B) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    B[col * 1024 + row] = A[row * 1024 + col];
}

正确写法（分块 + 共享内存）：

// 使用 32×32 分块，先合并读取 A 到共享内存
// 再从共享内存合并写入 B
__global__ void transpose_good(const float* A, float* B) {
    __shared__ float tile[32][33];  // 加 padding 避免 bank conflict
    int row = blockIdx.y * 32 + threadIdx.y;
    int col = blockIdx.x * 32 + threadIdx.x;
    
    // 合并读取：线程束访问 A[row][col..col+31] 连续地址
    tile[threadIdx.y][threadIdx.x] = A[row * 1024 + col];
    __syncthreads();
    
    // 合并写入：线程束访问 B[col][row..row+31] 连续地址
    int newRow = blockIdx.x * 32 + threadIdx.y;
    int newCol = blockIdx.y * 32 + threadIdx.x;
    B[newRow * 1024 + newCol] = tile[threadIdx.x][threadIdx.y];
}

误区与复盘：常见的内存访问陷阱

• 误区一：以为所有访问都会被硬件自动合并。实际上，只有当线程束内地址落在同一 128B 对齐区间时才会合并。若地址跨越两个缓存行边界，即使整体连续，也可能拆分为 2 个事务。例如线程束访问地址 120~248 字节，跨越 128B 边界，会触发 2 个事务。
• 误区二：忽略数据结构体的大小。当线程访问结构体数组（AoS，Array of Structures）时，若结构体大小不是 4/8/16 字节的整数倍，会导致地址对齐错乱。例如 struct { float x, y, z; } 大小为 12 字节，线程束访问 arr[i].x 时地址间隔 12B，无法合并。应改为 SoA（Structure of Arrays）布局。
• 误区三：过度依赖 L2 缓存“兜底”。虽然 L2 可以合并部分非连续访问，但 L2 带宽同样有限。非合并访问会占用 L2 的额外带宽和标签查找资源，即使命中缓存，延迟和功耗也更高。
• 误区四：忽略只读数据的纹理/只读路径。对于只读且访问模式不规则的数据，使用 __ldg() 或 const __restrict__ 可以走纹理缓存路径，提高缓存命中率，但这不能替代合并访问——它只是减少 L2 到 DRAM 的重复搬运。