3、内存访问模式：连续访问 vs 随机访问、对齐访问的重要性、合并访问技术
好，咱们接着聊内存带宽。上一节我讲了带宽的理论值和实际值之间的鸿沟，这节咱们就深入看看，到底是什么样的代码习惯，在悄悄偷走你的带宽。
说白了，GPU 内存的脾气很怪。你顺着它的毛捋，它给你满速跑。你要是逆着来，它直接给你降速到十分之一。我见过太多项目，算法算力明明够，就是被内存访问拖死的。
3.1 连续访问 vs 随机访问：天壤之别
先问个问题：你猜猜看，连续读 100 个 32 位整数，和随机读 100 个，速度能差多少？
我告诉你，在嵌入式 GPU 上，这个差距可以到 10 倍以上。为什么？
因为 GPU 的内存控制器是按「突发传输」工作的。你读地址 A，它会把 A 附近的一整块数据都搬进缓存。如果你接下来读 A+1，那数据已经在缓存里了，零等待。但如果你跳着读，缓存命中率暴跌，每次都要去显存里取，带宽自然就崩了。
连续访问：地址递增，像读数组一样 arr[0], arr[1], arr[2]...
随机访问：地址跳跃，像查哈希表一样 table[key1], table[key2]...
核心原则：让每个线程访问的地址尽量连续，且相邻线程访问相邻地址。
我在项目中遇到过这样一个案例：一个图像滤波算法，原本用指针链表存储邻域像素，随机访问严重，带宽利用率只有 12%。改成连续数组存储后，直接飙到 85%。改代码只花了两小时，性能翻了三倍。
3.2 对齐访问：别让 GPU 做额外劳动
对齐访问，听起来像个老生常谈的话题。但说实话，很多嵌入式开发者压根没注意过。
什么是内存对齐？简单说，就是你的数据地址，必须是它自身大小的整数倍。比如一个 4 字节的 int，地址最好是 4 的倍数；一个 16 字节的向量，地址最好是 16 的倍数。
为什么重要？因为 GPU 的内存总线宽度是固定的。以常见的 128 位总线为例，它一次就读 16 个字节。如果你的 int 放在地址 2 上，那 GPU 要读两次：先读 0-15，再读 16-31，然后拼出你要的 4 个字节。效率直接砍半。
| 数据类型 | 推荐对齐边界 | 未对齐惩罚 |
|---|---|---|
| int (4B) | 4 字节 | 2 次内存事务 |
| float4 (16B) | 16 字节 | 4 次内存事务 |
| double2 (16B) | 16 字节 | 4 次内存事务 |
我曾经踩过的坑：在结构体里随意排列成员，导致 GPU 端读取时大量未对齐访问。后来强制用 __attribute__((aligned(16))) 对齐，性能提升 40%。记住，结构体成员顺序也要按对齐要求排！
嗯，这里要注意：OpenCL 里用 cl_mem 分配的内存默认是对齐的，但如果你自己用指针做偏移，很容易破坏对齐。我建议你养成习惯，所有显存地址都按 16 字节对齐。
3.3 合并访问技术：让每个线程都「合群」
合并访问，这是 GPU 内存优化的核心中的核心。你想想看，GPU 有几百个核心，如果每个核心都单独去显存取数据，那内存控制器得忙死。所以 GPU 设计了一个机制：把多个线程的内存请求合并成一个大的请求。
具体来说，当一个 warp（32 个线程）同时访问内存时，硬件会检查这些地址是否连续。如果连续，就合并成一次或几次大的突发传输。如果不连续，就拆成很多小请求，效率极低。
合并访问的理想模式：
// 好的模式：线程 i 访问 data[i]
int tid = get_local_id(0);
float val = data[tid]; // 连续，可合并
// 坏的模式：线程 i 访问 data[permute[i]]
int tid = get_local_id(0);
float val = data[permute[tid]]; // 随机，不可合并我个人习惯在写 kernel 之前，先画一张「线程-地址映射图」。横轴是线程 ID，纵轴是访问地址。如果映射图是一条斜线，那就是完美合并。如果是一团散点，那就得重构数据布局了。
3.4 实战技巧：如何写出合并友好的代码
光讲理论没用，咱们来点实际的。我总结了三条铁律：
- 让线程 ID 对应连续地址：用
get_global_id(0)直接作为数组下标，这是最保险的做法。 - 矩阵访问要按行优先：如果你要遍历一个矩阵，让线程处理一行，而不是一列。因为行在内存中是连续存储的。
- 结构体数组 vs 数组结构体：尽量用 SoA（数组结构体）而不是 AoS（结构体数组）。
举个例子，假设你要处理粒子的位置和速度：
// AoS 模式（不推荐）
struct Particle {
float x, y, z;
float vx, vy, vz;
};
Particle particles[N];
// SoA 模式（推荐）
struct Particles {
float *x, *y, *z;
float *vx, *vy, *vz;
};
// 线程 i 访问 particles.x[i]，连续且可合并小技巧：如果你必须用 AoS，可以尝试用向量化加载。比如用 float4 一次读 4 个成员，虽然不如 SoA 完美，但比逐成员读取好得多。
我记得有一次优化一个 N-body 模拟，原本用 AoS，带宽利用率只有 30%。改成 SoA 后，利用率到了 90%，模拟规模直接翻倍。你想想看，同样的硬件，只是改了下数据布局，效果立竿见影。
3.5 避坑指南：常见的合并访问杀手
最后，我列几个我实际踩过的坑，你写代码时注意避开：
- 间接访问：用索引数组去取数据，比如
data[index[i]]。除非 index 是连续的，否则合并性极差。 - 条件分支导致的不连续：if 语句让部分线程跳过访问，剩下的线程地址不连续。可以考虑用掩码或重新组织数据。
- 跨步访问：比如每隔 4 个元素取一个，
data[i*4]。虽然连续，但步长太大，合并效率低。可以尝试让每个线程处理多个连续元素。
我曾经接手过一个图像处理项目，里面用了大量的查表操作，索引是随机生成的。结果带宽利用率不到 5%。后来我把表按访问模式重新排序，让相邻线程查相邻的表项，利用率直接拉到 70%。
好了，这节的内容就这些。总结一句话：让你的内存访问像排队一样整齐，别让 GPU 做无谓的散装搬运。下一节我会讲如何利用局部内存和共享内存来进一步优化，敬请期待。