4. 内存访问模式优化

各位同学，今天我们来聊聊CUDA性能优化里最核心、也最容易被忽视的一个话题——内存访问模式。说白了，就是你的线程怎么去读数据、写数据。我见过太多项目，计算逻辑写得漂亮，结果一跑起来，带宽利用率不到20%。为什么？内存访问没对齐，没合并。

4.1 什么是合并访问？

先问大家一个问题：GPU里有成百上千个线程同时跑，它们怎么从显存里拿数据？

答案是：以warp为单位。一个warp是32个线程，它们一起执行指令。当这32个线程要访问Global Memory时，硬件会尽量把这32次访问合并成一次或几次大的内存事务。

合并访问，就是指这32个线程访问的内存地址是连续的、对齐的。这样硬件就能用最少的事务数把数据搬进来。

我刚开始做CUDA时，总觉得这玩意儿玄乎。直到有一次，我写了一个矩阵转置的kernel，带宽只有理论值的10%。排查了半天，发现就是内存访问模式不对。从那以后，我每次写kernel前，都会先在纸上画一画线程和数据的对应关系。

4.2 对齐访问：别小看这128字节

对齐访问，说白了就是你的数据首地址得是128的倍数。为什么是128？因为GPU的L2 cache line是128字节。如果你地址不对齐，一次内存事务可能只用到一部分数据，剩下的带宽全浪费了。

举个例子：

// 不对齐的访问
float *data = (float*)malloc(N * sizeof(float));
// 假设data的地址是0x1004，不是128的倍数
// 那么访问data[0]时，硬件需要加载0x1000到0x107F的数据
// 但data[0]到data[31]只占用了0x1004到0x1084
// 这就跨了两个cache line，需要两次事务

// 对齐的访问
float *data_aligned;
cudaMalloc((void**)&data_aligned, N * sizeof(float));
// cudaMalloc保证地址至少对齐到256字节
// 访问data_aligned[0]时，一次事务搞定

4.3 合并访问的两种模式

合并访问主要分两种场景：顺序访问和跨步访问。我分别说说。

4.3.1 顺序访问：最理想的情况

线程0访问地址0，线程1访问地址1，线程2访问地址2……以此类推。这就是最完美的合并访问。

// 顺序访问示例：向量加法
__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];  // 完美合并！
    }
}

你看，每个线程访问自己的idx位置，地址连续。硬件会把32个线程的访问合并成一次128字节的事务。带宽利用率接近100%。

4.3.2 跨步访问：性能杀手

如果线程0访问地址0，线程1访问地址32，线程2访问地址64……这就麻烦了。每个线程之间隔了32个元素，地址不连续。

// 跨步访问示例：矩阵转置（错误写法）
__global__ void transpose_bad(float *in, float *out, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < width && col < width) {
        out[col * width + row] = in[row * width + col];  // 写操作跨步！
    }
}

这里写操作out[col * width + row]，同一个warp内的线程col是连续的，但col * width的跨度是width个float，也就是4*width字节。如果width很大，每个线程的地址相隔很远，硬件无法合并，只能每个线程单独发一次事务。性能直接崩盘。

4.4 如何最大化Global Memory带宽利用率？

说白了，就三点：

1. 保证合并访问：warp内线程访问连续地址
2. 保证对齐访问：首地址128字节对齐
3. 减少内存事务数：一次事务尽量用完128字节

我整理了一个表格，方便大家对照：

访问模式	warp内地址模式	内存事务数	带宽利用率
完美合并	连续，对齐	1	~100%
部分合并	连续，不对齐	2	~50%
跨步访问	间隔固定步长	32	~3%
随机访问	完全随机	32	~3%

嗯，这里要注意：跨步访问和随机访问，在极端情况下带宽利用率可能连1%都不到。所以写kernel时，一定要把「合并访问」刻在脑子里。

4.5 实战技巧：用共享内存打破跨步访问

遇到跨步访问怎么办？我的经验是：用共享内存做中转。先把数据从Global Memory以合并方式读到共享内存，再从共享内存以任意方式读取。

// 矩阵转置优化版：使用共享内存
__global__ void transpose_good(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    // 合并读：从Global Memory读到共享内存
    if (row < width && col < width) {
        tile[threadIdx.y][threadIdx.x] = in[row * width + col];
    }
    __syncthreads();
    
    // 合并写：从共享内存写到Global Memory
    int new_row = blockIdx.x * TILE_SIZE + threadIdx.y;
    int new_col = blockIdx.y * TILE_SIZE + threadIdx.x;
    if (new_row < width && new_col < width) {
        out[new_row * width + new_col] = tile[threadIdx.x][threadIdx.y];
    }
}

你看，读操作in[row * width + col]是合并的，写操作out[new_row * width + new_col]也是合并的。中间通过共享内存做了行列转置。这就是典型的「以空间换时间」。

4.6 知识体系总览

下面这张图，是我自己总结的内存访问优化知识体系。你看一眼，就能明白今天讲的所有内容是怎么串起来的。

这张图把今天的内容串起来了。你想想看，从Global Memory访问出发，三个分支——合并、对齐、带宽最大化，最终都指向同一个目标：让数据搬运的效率最高。而底部的实战技巧，就是我们在项目中真正要用到的手段。

4.7 小结

今天的内容就这些。总结一下：

• 合并访问是GPU内存优化的基石，warp内线程必须访问连续地址
• 对齐访问保证首地址是128的倍数，避免跨cache line
• 跨步访问是性能杀手，遇到就用共享内存或调整数据布局来破解
• 带宽利用率最大化，说白了就是让每次内存事务都物尽其用

嗯，这些内容看起来简单，但真正写代码时很容易忽略。我建议你每次写完kernel，都用nvcc的--ptxas-options=-v选项编译一下，看看内存事务数。如果事务数明显多于预期，那八成是访问模式出了问题。

好，今天就到这里。记住：内存访问模式对了，性能优化就成功了一半。