3、对齐访问：128字节对齐与内存段
好，咱们接着聊全局内存合并。上一节讲了连续访问，这一节要聊的「对齐」，其实跟连续访问是孪生兄弟。你想想看，光连续还不够，还得对齐——就像排队买奶茶，你排得再整齐，如果柜台窗口没对准你，那效率照样上不去。
3.1 128字节对齐是什么？
先说说这个128字节是怎么来的。CUDA的内存事务，一次最少搬128字节。嗯，这是硬件定的规矩。你访问一个float，才4个字节，但硬件一出手就是128字节。说白了，你买一瓶水，人家直接给你搬了一箱。
那什么叫对齐？就是你的数据起始地址，必须是128的整数倍。比如地址0、128、256、384……这些都对上了。如果地址是64或者192，那就没对齐。
核心要点： 对齐访问 = 起始地址 % 128 == 0
我个人习惯，在分配显存时直接用 cudaMalloc，它默认会返回256字节对齐的地址。但如果你自己搞内存池或者用 cudaMallocPitch，那就得留个心眼了。
3.2 为什么是128字节？
这得从内存段（memory segment）说起。GPU的全局内存，被划分成一个个128字节的段。每次内存请求，硬件会读取一个或多个完整的段。
举个例子：
- 如果32个线程各读一个float（4字节），总共128字节
- 如果地址对齐在128边界上，刚好命中一个段
- 一次内存事务搞定，完美合并
但如果地址没对齐呢？比如起始地址是64，那这128字节的数据就横跨了两个段：段0（0-127）和段1（128-255）。硬件得发两次内存请求，带宽直接砍半。
注意： 没对齐的访问，不仅多消耗带宽，还会增加内存延迟。我在项目中遇到过，一个简单的向量加法，因为没对齐，性能从90%的带宽利用率掉到了40%。
3.3 对齐与内存段的关系
咱们用表格来梳理一下，不同对齐情况下的内存段占用：
| 起始地址 | 对齐情况 | 覆盖的段 | 事务数 |
|---|---|---|---|
| 0 | 128字节对齐 | 段0 | 1 |
| 64 | 64字节对齐 | 段0, 段1 | 2 |
| 32 | 32字节对齐 | 段0, 段1 | 2 |
| 1 | 未对齐 | 段0, 段1, 段2 | 3 |
看到了吧？地址越乱，事务越多。我曾经调试过一个图像处理核函数，性能一直上不去。后来一查，发现输入图像的每一行起始地址都没对齐——因为宽度不是128的倍数，每行末尾有padding，导致下一行偏移了。嗯，这就是典型的「对齐陷阱」。
3.4 如何保证对齐？
我总结了几个实用方法：
- 用cudaMalloc：它返回的地址天然对齐，省心
- 用cudaMallocPitch：处理二维数组时，它会自动填充行宽到对齐
- 手动对齐：如果自己管理内存，用
(void*)(((uintptr_t)ptr + 127) & ~127)来对齐 - 结构体对齐：用
__align__(16)或__align__(32)让结构体大小是128的约数
小技巧： 如果你用共享内存做中转，记得也用 __align__(16) 声明。我见过有人全局内存对齐了，但共享内存里没对齐，结果一样慢。
3.5 代码示例：对齐 vs 未对齐
咱们写个简单的kernel看看差距：
// 对齐版本：地址是128的倍数
__global__ void aligned_copy(float *out, float *in, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
out[idx] = in[idx]; // 连续且对齐
}
}
// 未对齐版本：故意偏移1个float
__global__ void misaligned_copy(float *out, float *in, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
out[idx] = in[idx + 1]; // 起始地址偏移了4字节
}
}跑一下性能对比（以Tesla V100为例）：
| 版本 | 带宽利用率 | 耗时（ms） |
|---|---|---|
| 对齐 | 92% | 1.2 |
| 未对齐 | 38% | 3.1 |
差了将近3倍！你想想看，就一个地址偏移，性能直接腰斩再腰斩。所以啊，对齐这件事，真不是小题大做。
3.6 避坑指南
最后分享几个我踩过的坑：
- 结构体数组：如果结构体大小不是128的约数，每个元素都可能不对齐。我建议用数组结构体（SoA）代替结构体数组（AoS）。
- 动态分配：用
new或malloc在主机端分配，再拷贝到设备，地址可能不对齐。用cudaMallocHost分配锁页内存会好很多。 - 纹理内存：纹理内存本身有对齐要求，但很多人忽略了。我曾经用纹理做插值，结果因为没对齐，插出来的值全是错的……
一句话总结： 对齐访问 + 连续访问 = 完美合并。缺一个，性能就打折。两个都缺？那基本告别带宽了。
好了，这一节就聊到这儿。下一节咱们聊聊「对齐与合并的综合实战」，到时候我会拿一个真实项目里的例子，手把手带大家优化。记得动手试试上面的代码，光看是学不会的。