2、合并访问的条件：对齐、连续、线程束粒度

好，咱们接着聊。上一节我讲了合并访问为什么重要，说白了就是让显存带宽别浪费。那这一节，我们来拆解一下：到底满足什么条件，才能让全局内存访问真正“合并”？

我个人习惯把条件归纳为三个关键词：对齐、连续、线程束粒度。缺一个，性能就可能打折扣。

2.1 对齐：起始地址要“规矩”

什么叫对齐？简单说，就是你的数据起始地址，必须是某个对齐单位的整数倍。

对于全局内存合并访问，硬件要求是：一个线程束（32个线程）访问的起始地址，必须对齐到32字节。为什么是32字节？因为一个缓存行（cache line）通常是128字节，而一次合并访问的最小粒度是32字节。嗯，这里要注意，如果你地址没对齐，硬件会帮你拆成多次访问，性能直接腰斩。

我在项目中遇到过这样一个坑：有个同事写了一个kernel，数据从某个偏移量开始读，结果性能比预期慢了一倍。我一看代码，发现他用了 float* data + 1 作为起始地址。你想想看，float是4字节，偏移1个就是4字节，4 % 32 != 0，地址没对齐。改成 float* data + 8（偏移32字节），性能立刻恢复正常。

2.2 连续：线程访问的地址要“挨着”

对齐只是第一步。第二个条件是：线程束内32个线程访问的地址，必须是连续的。

什么叫连续？就是线程0访问地址A，线程1访问地址A+1，线程2访问地址A+2……以此类推。说白了，就是地址按线程ID递增，中间不能跳。

为什么会这样？因为硬件一次可以读取连续的一段内存，比如128字节。如果地址是连续的，一个内存事务就能搞定。如果地址是跳跃的，比如线程0读A，线程1读A+100，那硬件就得发起多次事务，带宽利用率直线下降。

我举个例子，假设你有一个float数组，每个线程读一个元素：

// 连续访问（合并）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];

// 非连续访问（不合并）
int idx = blockIdx.x * blockDim.x + threadIdx.x * 2;
float val = data[idx];

你看，第一个例子中，线程0读data[0]，线程1读data[1]，地址连续，完美合并。第二个例子中，线程0读data[0]，线程1读data[2]，中间跳了一个，地址不连续，性能就差了。

2.3 线程束粒度：以32线程为单位

第三个条件，其实是对前两个条件的补充：合并访问是以线程束（warp）为单位的。

你想想看，GPU调度执行的最小单位是线程束，不是单个线程。所以硬件在发起内存访问时，也是以线程束为单位来检查是否满足合并条件。也就是说，对齐和连续，都是针对一个线程束内的32个线程而言的。

举个例子，假设你有一个kernel，每个block有128个线程。硬件会把这128个线程分成4个线程束（warp 0~3）。每个线程束独立判断自己的访问是否合并。warp 0的32个线程地址连续且对齐，那它就能合并；warp 1的地址不连续，那它就不能合并。两个互不影响。

我记得有一次优化一个矩阵转置的kernel，发现性能始终上不去。后来我用nvprof一分析，发现每个线程束的访问模式不一样：有的warp是连续读，有的warp是跳跃读。原因是我把线程块维度设成了16x8，导致线程ID映射到地址时，部分warp的地址不连续。改成32x4后，所有warp都连续了，性能提升了30%。

2.4 三个条件的关系

我把这三个条件总结一下，方便你记忆：

条件	含义	违反后果
对齐	起始地址是32字节的整数倍	一次访问拆成多次，带宽减半
连续	线程束内地址按线程ID递增	多次内存事务，带宽利用率低
线程束粒度	以32线程为单位检查合并	部分warp合并，部分不合并，整体性能差

这三个条件其实是层层递进的。对齐是基础，连续是关键，线程束粒度是视角。你写代码时，只要记住一句话：让每个线程束的32个线程，访问一段连续且对齐的内存区域。做到这一点，合并访问就稳了。

好，这一节就到这里。下一节我会讲一个实际案例——矩阵乘法的合并访问优化，到时候你会看到这些条件是怎么落地的。