全局内存访问模式：合并访问与对齐要求

各位同学，今天我们来聊聊全局内存访问。说实话，这是GPU性能优化里最基础也最关键的一环。我见过太多项目，核函数写得花里胡哨，结果一跑性能惨不忍睹——十有八九都是内存访问模式出了问题。

什么是合并访问？

先问大家一个问题：GPU里有成百上千个线程同时访问全局内存，你觉得它们是怎么读取数据的？

嗯，不是每个线程独立发请求。实际上，GPU会把32个线程（一个warp）的内存请求合并成尽可能少的几次事务。这就是合并访问（Coalesced Access）的核心思想。

合并访问：当一个warp内的所有线程访问连续的内存地址时，硬件可以将这些请求合并成少数几个大粒度的内存事务。说白了，就是让数据像排队一样整整齐齐地进来。

反过来，如果线程访问的地址是分散的，硬件就得拆成很多小事务，效率直线下降。我在项目中遇到过这样的情况：同样的计算量，合并访问比非合并访问快了将近10倍。这不是夸张，是真的。

对齐要求：别让你的数据歪着放

合并访问还有一个好兄弟，叫对齐。什么意思呢？就是内存地址得是某个粒度的整数倍。

举个例子，假设你有一个float数组，每个float占4字节。如果数组的首地址是64字节对齐的，那么第一个warp访问前32个float（128字节）时，硬件只需要一次事务就能搞定。

但如果地址没对齐，比如偏移了16字节，那硬件可能就得拆成两次事务。一次读前半段，一次读后半段。你想想看，这多出来的开销是不是很冤枉？

对齐粒度	事务大小	最佳实践
32字节	32字节事务	适用于char/short类型
64字节	64字节事务	适用于int/float类型
128字节	128字节事务	适用于double/vector类型

我个人习惯在分配内存时，直接用cudaMalloc或hipMalloc，它们默认会返回至少256字节对齐的地址。但如果你自己搞内存池或者用共享内存，就得小心了。

小技巧：用 cudaMallocPitch 或 hipMallocPitch 分配二维数组，可以自动保证每行的对齐。我在做图像处理时经常用这个，省心不少。

Stride访问：一个常见的坑

Stride访问，说白了就是线程访问内存时跳着走。比如每个线程读一个float，但间隔是8个float。这种情况在矩阵转置、图像处理里特别常见。

为什么会这样？我们来看一段代码：

// 非合并访问：按列读取
__global__ void read_by_column(float* data, int width) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    float val = data[row * width + col];  // 这里没问题
}

// 合并访问：按行读取
__global__ void read_by_row(float* data, int width) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    float val = data[row * width + col];  // 这里有问题！
}

注意看第二个核函数。线程索引的分配方式变了，导致同一个warp内的线程访问的地址间隔了width个float。如果width是1024，那间隔就是4096字节。硬件根本没法合并，只能拆成32次独立事务。

我曾经接手过一个流体模拟的项目，里面有个核函数跑得特别慢。一查，就是这种Stride访问。改成按行读取后，性能直接翻了6倍。嗯，有时候优化就是这么简单粗暴。

Stride访问的优化策略

遇到Stride访问怎么办？我有几个常用的招数：

改变数据布局：把AoS（数组结构体）改成SoA（结构体数组）。这是最根本的解法。
使用共享内存做转置：先让每个线程把数据读到共享内存里，再按需要的顺序写回全局内存。
调整线程块大小：有时候把blockDim.x设成warp大小（32）的倍数，能缓解一些问题。
向量化加载：用float4或int4一次读多个数据，减少事务次数。

举个例子，用共享内存做转置：

__global__ void transpose_coalesced(float* in, float* out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    // 合并读取到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 转置后合并写回
    out[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

你看，读和写都是合并的，中间只多了一次共享内存的转置。这个技巧我在做矩阵运算时用了无数次。

注意：共享内存的大小是有限的。如果TILE_SIZE设得太大，会导致每个SM上活跃的线程块数量减少，反而影响性能。我一般建议TILE_SIZE取32或16，具体要看你的GPU型号。

知识体系总览

下面这张图总结了全局内存访问模式的核心逻辑：

说实话，全局内存访问优化是GPU编程的基本功。你把这个搞明白了，后面学共享内存、寄存器优化都会轻松很多。我建议你在写核函数之前，先花30秒想想：这个warp里的32个线程，它们访问的地址是连续的吗？对齐了吗？

嗯，养成这个习惯，你的核函数性能不会差到哪里去。

避坑指南：我曾经在写一个粒子模拟的核函数时，为了图方便把粒子的位置和速度放在一个结构体数组里（AoS）。结果每个warp访问粒子数据时，位置和速度交错排列，导致完全无法合并。改成SoA后，性能提升了8倍。从那以后，我写数据结构都会先想想GPU的访问模式。

好了，关于全局内存访问模式就讲到这里。记住三个关键词：合并、对齐、连续。下次写核函数的时候，多想想这三个词。

📚 课程目录

全局内存访问模式：合并访问与对齐要求

什么是合并访问？

对齐要求：别让你的数据歪着放

Stride访问：一个常见的坑

Stride访问的优化策略

知识体系总览