在高性能计算的世界里，尤其是当你一头扎进GPU编程的深水区，有一个词你绝对绕不开——内存合并（Memory Coalescing）。别被这听起来高大上的术语吓到，它其实没那么玄乎，但一旦你理解并用对了，你的程序性能可能直接起飞！今天，我们就用最接地气的方式，把内存合并这件事掰开揉碎，讲得明明白白。
首先，咱们得知道GPU的内存架构和CPU不太一样。在英伟达的CUDA架构中，全局内存（Global Memory）是直接连到显卡上的物理显存，比如GDDR6或者HBM这类高速显存。这些显存带宽超高，动辄几百GB每秒，听起来是不是很爽？但问题来了——它们的访问延迟特别高！比你家电脑用的DDR5内存还要慢不少。为啥？因为显存用的是DRAM技术，靠电容充放电来存数据，而电容充放电的速度受物理限制，比如发热、功耗、芯片面积等等，根本快不起来。
那问题就来了：如果每个线程都单独去读一个内存地址，GPU的内存带宽根本跑不满，大量时间都浪费在等数据上，性能直接打骨折。这时候，内存合并就闪亮登场了！
内存合并本质上是一种硬件优化技术，它的核心思想特别简单：与其让32个线程各自发32次内存请求，不如把这32次请求“打包”成一次物理访问，一次性把需要的数据全拉回来。这样，不仅减少了访问次数，还充分利用了DRAM的内部机制。
说到DRAM，这里得插一句：当你访问一个内存地址时，DRAM其实不是只读那一个地址，而是会顺手把后面一整块连续的数据一起读出来，这个操作叫“突发传输”（Burst Transfer）。比如，在现代GPU上，一次突发传输通常是128字节。巧的是，一个CUDA线程束（Warp）正好有32个线程。如果每个线程读一个32位的浮点数（也就是4字节），32×4=128字节——完美匹配！只要这32个线程读的地址是连续的、对齐的，硬件就能用一次突发传输搞定全部数据，这就是最理想的内存合并状态。
反过来说，如果你的线程读的地址东一个西一个，七零八落，那GPU就不得不发起多次突发传输，甚至可能一次只服务一个线程，其他31个干等着。这时候，带宽利用率暴跌，程序慢得像蜗牛。
为了让大家更直观感受内存合并的威力，我们来看一个经典实验。假设我们写一个CUDA核函数，让每个线程从一个大数组里读数据，但读的时候不是挨着读，而是隔几个位置读一次——这个间隔就叫“步长”（Stride）。当步长为1时，所有线程读的是连续地址，内存合并效果拉满；但当步长变成2、4、8……情况就急转直下。
比如在特斯拉T4显卡上跑这个实验，数组大小256MB，结果非常震撼：步长为1时，内存带宽能达到206 GB/s；步长变成2，直接掉到130.5；步长4，只剩68.8；到了步长8，只剩33.8。基本上每翻一倍步长，带宽就砍一半！这说明什么？说明内存合并失效后，性能断崖式下跌。
但有趣的是，当步长超过16之后，带宽下降速度变缓了，比如步长32、64、128时，带宽从15.2慢慢降到11.2。这背后其实还有别的因素在起作用，比如GPU内部的地址转换缓存（TLB）开始频繁失效，或者缓存局部性彻底崩坏，导致整个内存子系统效率进一步恶化。
说到这里，你可能会问：CPU难道没有类似机制吗？当然有！CPU也有缓存行（Cache Line），通常是64字节，也会把连续内存一起加载。但关键区别在于——CPU的这些优化对程序员是透明的，你不用管，硬件自动搞定。而GPU不一样！在CUDA编程中，内存访问模式完全由你控制。硬件不会帮你“猜”你要读什么，所以你必须主动写出能被合并的代码。这也是为什么GPU编程既强大又“痛苦”——性能上限极高，但下限也极低，全看你怎么写。
那么，怎么写出能内存合并的代码呢？核心原则就一条：让同一个线程束（Warp）里的32个线程，访问连续且对齐的内存地址。最简单的做法就是让线程ID直接对应数组下标。比如threadIdx.x = 0 的线程读 in[0]，threadIdx.x = 1 的读 in[1]，以此类推。这样，一个Warp读的就是 in[0] 到 in[31]，完美连续，128字节一次搞定。
但现实中，数据结构往往更复杂。比如你处理的是二维图像，或者结构体数组（AoS），这时候就容易踩坑。比如每个结构体包含x、y、z三个float，如果你让每个线程读一个结构体的x分量，那实际内存地址间隔是12字节（3×4），32个线程的访问地址就不是连续的，内存合并就失效了。这时候，高手的做法是改用“结构体转数组”（SoA）布局——把所有x放一块，所有y放一块，所有z放一块。这样，读x的时候就是连续内存，合并效果拉满。
顺便提一句，这篇文章的核心思想其实源自一本经典教材——《大规模并行处理器编程》（Programming Massively Parallel Processors）第四版的第6.1节。这本书的作者之一是David B. Kirk，他曾是英伟达的首席科学家，也是CUDA架构的奠基人之一。另一位作者Wen-mei W. Hwu是伊利诺伊大学香槟分校的教授，在并行计算领域享有盛誉。他们写的这本书，堪称GPU编程的“圣经”。而文中还提到Ulrich Drepper的《每个程序员都应该了解的内存知识》，这篇长文在系统编程圈子里也是神作，深入剖析了从CPU缓存到内存控制器的底层机制。
回到实战，如果你正在写CUDA程序，一定要时刻问自己：我的内存访问模式合并了吗？可以用Nsight Compute这样的性能分析工具看看“内存吞吐效率”或者“L2缓存命中率”。如果发现带宽远低于理论峰值，十有八九是内存访问出了问题。
最后再强调一遍：在GPU的世界里，计算不是瓶颈，内存才是！你的浮点运算再快，数据读不进来也是白搭。而内存合并，就是打通这条“数据高速公路”的关键钥匙。掌握它，你就能让GPU真正发挥出恐怖的并行吞吐能力。
所以，别再让你的GPU“饿着肚子干活”了！从今天起，检查你的访存模式，拥抱内存合并，让你的CUDA程序跑出火箭速度！
（本文技术内容综合自英伟达开发者博客、经典教材及系统性能研究文献，旨在为GPU开发者提供实用指导。）