3. 全局内存合并访问：对齐与合并条件、Stride访问与Gap访问的性能影响、Padding技术

说到全局内存访问，我估计不少同学都踩过坑。明明算法设计得挺漂亮，一跑起来性能就是上不去。这时候十有八九是内存访问模式出了问题。今天咱们就聊聊全局内存的合并访问——说白了，就是怎么让GPU一次性多搬点数据，别让带宽白白浪费。

3.1 合并访问的基本概念

先问个问题：为什么GPU要搞合并访问？

你想想看，GPU有几千个核心，但显存带宽就那么点。如果每个线程都单独去显存里取数据，那带宽根本不够用。所以NVIDIA设计了一个机制：当连续线程访问连续地址时，硬件会把多次小请求合并成一次大请求。这就是合并访问。

核心要点：合并访问能让显存带宽利用率接近理论峰值。不合并的话，可能连10%都用不到。

我记得刚做CUDA开发那会儿，写了个简单的向量加法，结果带宽只有理论值的5%。当时百思不得其解，后来才发现是内存对齐出了问题。嗯，这事儿后面细说。

3.2 对齐条件

合并访问有两个前提：对齐和连续。

对齐指的是内存地址必须是32字节、64字节或128字节的整数倍。为什么是这些数？因为GPU一次事务（transaction）能搬32字节、64字节或128字节。

看个例子：

// 对齐访问（好）
float* data;
data = (float*)cudaMalloc(&data, N * sizeof(float));
// cudaMalloc默认返回256字节对齐的地址

// 非对齐访问（差）
float* misaligned = data + 1;  // 偏移了4字节
// 这会导致一次事务要读两个cache line

避坑指南：我曾经在项目中用cudaMallocPitch分配二维数组，结果忘了考虑pitch的对齐要求，导致后续的kernel性能暴跌。后来查了半天才发现是pitch没对齐到128字节。

3.3 合并条件

对齐只是第一步。合并访问还需要满足：线程ID和内存地址一一对应。

具体来说：

线程0访问地址0
线程1访问地址1
线程2访问地址2
……

这样硬件就能把32个线程的请求合并成一次128字节的事务。

看代码对比：

// 合并访问（好）
__global__ void coalesced_kernel(float* data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        float val = data[idx];  // 连续线程访问连续地址
        // 处理数据...
    }
}

// 非合并访问（差）
__global__ void non_coalesced_kernel(float* data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        float val = data[idx * 2];  // 间隔访问，无法合并
        // 处理数据...
    }
}

3.4 Stride访问与Gap访问

Stride访问就是每隔固定步长取数据。Gap访问则是中间有空洞。这两种模式都会破坏合并访问。

Stride访问的性能影响：

Stride大小	带宽利用率	说明
1（连续）	~100%	完美合并
2	~50%	一半带宽浪费
4	~25%	四分之三浪费
随机	<10%	基本没有合并

为什么会这样？因为Stride=2时，每个线程只用到事务中一半的数据。另一半被浪费了。Stride越大，浪费越严重。

个人经验：我优化过一个矩阵转置的kernel。原始版本用Stride访问，带宽只有30GB/s。改成共享内存做转置后，带宽直接飙到200GB/s。差距就是这么明显。

3.5 Padding技术

Padding说白了就是加一些无用数据，让内存访问模式变得连续。

最常见的场景是二维数组。假设有个矩阵M[N][M]，按列访问时：

// 按列访问（差）
for (int i = 0; i < N; i++) {
    float val = matrix[i * M + col];  // 跨行访问，Stride = M
}

// 加Padding后（好）
int padded_M = M + padding;  // 让padded_M是32的倍数
for (int i = 0; i < N; i++) {
    float val = padded_matrix[i * padded_M + col];  // 连续访问
}

Padding的原理很简单：让每行的起始地址对齐到cache line边界。这样即使按列访问，也能保证每行内的数据是连续的。

实战建议：Padding的大小通常是32字节的整数倍。我个人习惯用128字节，因为这是L2 cache line的大小。多花点显存，换来带宽翻倍，这笔买卖划算。

3.6 知识体系图

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

3.7 实战经验总结

说了这么多，最后分享几个实用技巧：

优先保证连续访问：能用连续访问就别用Stride。实在不行就考虑转置数据。
对齐检查不能省：每次分配内存后，用cudaPointerGetAttributes检查对齐情况。
Padding要舍得：多花10%的显存，换来2-3倍的带宽提升，这笔账怎么算都划算。
用profiler验证：ncu或nvprof都能显示合并访问情况。我习惯先跑一遍profiler，看看global load/store效率。

最后提醒一句：别以为写了连续访问就万事大吉。我记得有一次，代码看起来是连续的，但因为结构体大小不是32的倍数，导致每个线程实际访问的地址有偏移。这种坑防不胜防，一定要用profiler验证。

好了，全局内存合并访问就聊到这儿。记住一句话：让连续线程访问连续地址，让地址对齐到cache line。做到这两点，你的kernel性能至少不会太差。

想要更多资料，联系可可

📚 课程目录

3. 全局内存合并访问：对齐与合并条件、Stride访问与Gap访问的性能影响、Padding技术

3.1 合并访问的基本概念

3.2 对齐条件

3.3 合并条件

3.4 Stride访问与Gap访问

3.5 Padding技术

3.6 知识体系图

3.7 实战经验总结