让内存自己干活：CXL近数据计算如何把数据库快成闪电！
CPU访问CXL内存虽然快，但频繁跑腿还是累。本文提出M²NDP方案，让CXL内存自己干活。它把函数调用伪装成普通内存访问，实现微秒级任务派发；再用轻量级“微线程”把内存带宽榨干。最终，数据库、推荐系统等任务最高快128倍，省电88%。
期刊/发表日期/原文标题/作者背景
原文标题：Low-overhead General-purpose Near-Data Processing in CXL Memory Expanders
作者：Hyungkyu Ham, Jeongmin Hong 等（韩国POSTECH大学、首尔国立大学、SK海力士）
发表日期：2024年4月（arXiv:2404.19381v2）
作者背景：学术界与工业界联合团队，专注计算机体系结构、内存系统、CXL互连技术
传统CXL内存就是个大仓库，但CPU来回跑腿累成狗
咱们先搞清楚CXL是啥。Compute Express Link，说白了就是一种高速线缆，能让CPU连上外面的大内存条。这玩意儿现在很火，因为它比以前的PCIe快得多。CPU发一条“把地址0x1234的数据给我”的命令，大概150纳秒就能拿到数据。这个速度跟CPU访问隔壁插槽的内存差不多，完全可以接受。
所以大家就开始搞CXL内存扩展。以前你服务器只有256GB本地内存，现在插个CXL盒子，里面塞2TB内存，成本还低。数据库、推荐系统、大语言模型这些吃内存的大户，一下子就宽敞了。
但是问题来了。CXL那条线的带宽有限。本地内存内部带宽可能有几百GB每秒，但CXL链路只有几十GB。这就好比你家水龙头出水很猛，但接了一根细管子通到院子里，院子那边水就小了。
论文里有个图特别逗：一个内存密集型程序，数据放本地内存跑得飞快，放CXL内存里直接慢了将近10倍。为啥？因为程序不停地在内存里搬数据，CXL那条细管子堵死了。还有一类程序对延迟敏感，比如键值存储，每次查找都要等CXL回复，尾延迟暴涨。
所以聪明的CPU就被迫当跑腿的了。它不停地发请求：“给我这个地址的数据”、“把这个值加1”、“把结果存回去”。每次都要走CXL链路，CPU自己累不说，链路还堵。
这时候有人就想了：让CXL内存自己干活不就行了？数据就在那儿，你直接在内存旁边把计算做了，只把最后结果告诉CPU。这就是近数据计算。
之前的近数据计算方案要么挑食，要么太贵，要么沟通慢得像寄信
有些前辈试过做专门的近数据计算单元。比如只干推荐系统的活，或者只干基因序列比对的活。效果确实好，但换个任务就抓瞎。要是每个应用都做一个专用芯片，那成本上天了。FPGA倒是能改，但写FPGA程序跟受刑似的，一般人搞不定。
另一些人干脆把CPU核或者GPU核塞进CXL内存里。这就像给仓库配了个全能管家。但是CPU核那套复杂的乱序执行逻辑、大缓存，在内存旁边根本用不上。内存密集型任务本来就没什么计算量，大部分时间在等数据。你配一个超标量乱序执行的核心，纯属大炮打蚊子。而且CPU核贵啊，面积大、功耗高。GPU核呢？一个SM（流多处理器）要管32个线程捆成一捆，如果这32个线程走的路不一样（分支发散），整个一捆就慢下来。内存密集型任务里这种情况很常见。
更麻烦的是沟通方式。以前CPU怎么让设备干活？用CXL.io协议，走PCIe那套。你要启动一个任务，CPU得在内核态写命令，设备那边有个环形缓冲区，两边还要不断地更新指针。一来一回好几微秒就过去了。对于那种运行时间很短的任务（比如几微秒），光沟通时间就占了大部分。这就像你打电话点外卖，外卖还没做，你跟接线员聊了半小时。
还有更原始的办法：用MMIO直接写设备寄存器。这个快一点，但寄存器就那么几个，多个程序没法同时用，还得切内核态。
核心思想：把函数调用伪装成普通内存读写，让沟通快如闪电
这篇论文的脑洞很简单：CXL.mem协议本来就很快，那咱们把“启动任务”、“查询状态”这些命令，伪装成普通的内存读写包。CPU发一个“写”操作到某个特殊地址，CXL内存一看这个地址，就知道你要启动哪个函数。再发一个“读”操作到同一个地址，就能拿到返回值。
具体做法是这样：在CXL内存里划出一块区域，叫M²func区。每个进程都有自己的这块区域。CXL内存的入口处放一个包过滤器，专门看地址。如果是普通地址，就正常读写内存。如果是M²func区里的地址，就触发相应的函数。
比如说你要注册一个核函数。你的程序执行一个store指令，把核函数的代码地址、需要多少寄存器这些信息写进去。这个写操作会被过滤，交给CXL控制器。控制器就把这个核函数登记在案，返回一个ID。这个ID怎么拿？你再执行一个load指令，读同样的地址，控制器就把返回值塞给你。
就这么简单。整个过程不切内核态，不走复杂的协议栈，就是两三个内存访问的功夫。CXL.mem的往返延迟大概70到150纳秒。所以一个任务启动，几百纳秒就搞定。相比之下，传统方案要几微秒甚至十几微秒。
你看图5那个时间线对比。同样一个短任务，M²func的总时间比环形缓冲区方案少了三分之一到四分之三。对于那种只跑几微秒的任务，这简直就是救星。
而且这个方案不用改CXL标准，也不用改CPU硬件。包过滤器是CXL内存自己加的，CPU那边完全不知道你在耍花招。它还以为自己在读写普通内存呢。兼容性满分。
微线程设计：让内存带宽被榨干，同时省掉地址计算的苦力
任务派发快了，接下来就要让干活的人也快。CXL内存里的计算单元得能同时处理海量请求，把内存带宽跑满。
传统CPU靠乱序执行来隐藏内存延迟，但逻辑太复杂。GPU靠细粒度多线程：一个核心上跑几千个线程，这个线程等内存，就换下一个。但GPU有两大毛病。
毛病一：地址计算啰嗦。在GPU里，每个线程要知道自己是第几个块里的第几个线程，然后用这些编号算出自己要处理哪个数据。代码里全是threadIdx.x, blockIdx.x, blockDim.x这些玩意。算一个地址要好几条指令。内存密集型任务本身就没什么指令，地址计算能占不小比例。
毛病二：线程创建太粗。GPU一次创建一个线程块，里面几十到几百个线程。块里有的线程干完活了，但块还没完，它占着的寄存器等资源就不能释放。造成浪费。
这篇论文提出的M²μthr对症下药。
首先，微线程的创建跟内存地址直接挂钩。你启动一个核函数的时候，指定一块地址范围作为“线程池”。每个微线程自动跟这块地址里的一个位置绑定。比如地址范围是0x1000到0x2000，每个微线程处理32字节，那么第一个微线程处理0x1000到0x101F，第二个处理0x1020到0x103F，以此类推。微线程启动时，寄存器x1放地址，x2放偏移量。核函数里直接用这些值去访存，根本不用算。论文测下来，静态指令数减少了3%到18%。
其次，微线程是单独创建的，不是成块创建。一个微线程干完了，立马释放资源，下一个就能顶上。资源利用率高得多。论文里一个图显示，对于某个图算法，GPU的活跃线程比例只有一半多，而M²NDP达到了九成以上。
另外，这套设计用RISC-V向量扩展。标量指令处理循环变量、地址，向量指令处理大批量数据。避免了GPU那种全SIMT的冗余。
片上还有一块共享内存。所有运行在同一个NDP单元上的微线程都能共享它，不像GPU的共享内存只能在一个线程块内共享。这样减少了很多全局内存访问。
整个干活流程：从启动到执行到结束，每一步都优化
现在把两个部分串起来看一个完整例子。
假设你要做向量加法。向量A、B、C都放在CXL内存里，每个向量有100万个元素。
第一步，注册核函数。你的程序调用ndpRegisterKernel，把向量加法的二进制代码位置、它需要几个整数寄存器、几个向量寄存器告诉CXL内存。M²func机制用一个写操作完成这个注册。CXL控制器把这个核函数存起来，返回一个ID，比如42。
第二步，启动核函数。调用ndpLaunchKernel，参数包括核函数ID、线程池区域（就是向量C的地址范围）、同步方式（异步还是同步）。又是一个写操作。CXL控制器收到后，通知微线程生成器：来活了，给向量C里的每个32字节块创建一个微线程。
微线程生成器开始干活。它给每个微线程分配一个槽位，里面有程序计数器、寄存器基址。每个微线程的x1寄存器被设为对应的地址（比如C[0]的地址），x2设为偏移量0。然后微线程开始执行核函数代码。代码里大致是：用x1的地址加载A的元素和B的元素，用向量加法指令加在一起，存回x1的地址。
如果用的是同步启动，CPU那边会发一个读请求到M²func区等待，直到所有微线程执行完。如果是异步启动，写请求直接返回，CPU接着干别的事，后面再调用ndpPollKernelStatus来查进度。
假设CXL内存里有32个NDP单元，每个单元有4个子核心，每个子核心能同时跑16个微线程。那么总共能同时跑2048个微线程。每个微线程处理32字节，一次就能处理64KB数据。内存带宽很容易就跑满了。
论文里测了一个数据库过滤操作。原来的CPU方案，数据要从CXL内存读到CPU，CPU判断条件，再写回去。用了M²NDP后，直接在CXL内存里判断，只把结果位图传回CPU。速度提升了128倍，CXL内存的内部带宽利用率达到90%以上。
实验结果：不仅快，还省电，而且面积小得惊人
论文用了一个周期精确的模拟器，模拟了完整的CXL内存和NDP单元。对比了几种方案：
- 基线：普通CPU或GPU，配上被动CXL内存
- CPU-NDP：把高性能CPU核塞进CXL内存
- GPU-NDP：把GPU SM塞进去，分等性能、等面积、4倍性能、16倍性能几种配置
测试的工作负载很全：数据库OLAP查询、键值存储、图算法、推荐系统的嵌入表查找、大语言模型推理。
结果如下：
对于OLAP，M²NDP比CPU基线快73倍（最高128倍）。比用32个高性能CPU核的CPU-NDP还快34%。因为那些CPU核虽然强，但它们的缓存和乱序逻辑在这个场景下纯属浪费。
对于键值存储，M²NDP用异步启动和状态轮询，把尾延迟降了79%。关键是能同时处理多个请求，吞吐量比寄存器方案高了47倍。
对于GPU任务，M²NDP比GPU基线快6.35倍。比等面积的GPU-NDP（只有16个SM）快1.41倍，比等性能的GPU-NDP（8个SM）快5.48倍。甚至比128个SM的GPU-NDP（16倍性能）还快24%。因为SM太多反而把内存行缓冲局域性破坏了。
M²func本身贡献多大？对比用CXL.io环形缓冲区的方案，M²func额外提速了23.8%。对于极短的任务（比如小批量DLRM），提速能达到2.4倍。
能耗方面，M²NDP比CPU基线省电84%，比GPU基线省电78%。因为它在CXL内存旁边干活，数据不用长途跋涉。
硬件成本呢？32个NDP单元总面积才26.4平方毫米（7纳米工艺）。这大概相当于一个指甲盖的几分之一。相比之下，一个高性能CPU核心就要几十平方毫米。
扩展到更多设备和开关：一个NDP单元不够，那就堆一堆
如果CXL内存不止一个，可以连好几个。CXL 3.0支持设备之间直接访问（P2P）。你可以在多个CXL内存上做模型并行：把大语言模型的一层放一个内存，推理时数据串行流过。论文测试了用8个CXL-M²NDP跑OPT-30B模型，加速比接近线性（7.69倍）。因为每个设备只处理自己那部分数据，P2P访问开销被分摊了。
还有一种玩法：把NDP单元做在CXL开关里。开关连着好几个被动CXL内存条。开关里的NDP单元可以访问所有这些内存条的数据。这样你可以单独增加内存容量而不增加计算能力。比如你已经有32个NDP单元，觉得够了，但还想加内存，就挂几个被动内存条。开关里的NDP单元照样能处理它们。测试下来，连8个被动内存条，性能提升7倍左右。
对比专用芯片：通用方案没输多少，但灵活性强太多
论文还跟几个专用近数据计算方案做了对比。比如专门做近似最近邻搜索的CXL-ANNS，专门做推荐系统的RecNMP，专门做基因组分析的Beacon。
结果M²NDP的性能跟这些专用方案差不多，平均只差6.5%。因为内存密集型任务的瓶颈是内存带宽，不是计算。只要能把带宽用满，通用核心和专用核心的区别不大。M²NDP用轻量级微线程做到了81.6%的内存带宽利用率，专用核心可能高几个百分点，但差距很小。
而专用核心的代价是：换个任务就歇菜。M²NDP跑完推荐系统，转身就能跑图算法，还能跑数据库查询。灵活性碾压。
写在最后：内存计算的时代真要来了
这篇论文的精髓就两句话：用内存访问来模拟函数调用，把通信开销降到几百纳秒；用轻量级微线程榨干内存带宽，同时省掉地址计算的破事。
未来的CXL内存很可能不再是个傻大个存储盒子，而是一个能自己干活的智能助手。你往某个地址写几个字节，它就开始算；你从那个地址读，它就给你结果。数据库的谓词下推、图算法的迭代计算、推荐系统的嵌入表查找，全能在内存里完成。
论文里甚至展望了更酷的场景：你写一个正则表达式，内存自己去扫描整个数据集，只把匹配的行返回给你。CPU彻底从数据搬运工变成真正的决策者。
当然，路还长。RISC-V的软件生态还没那么完善，编译器要改，库要重写。
X上七嘴八舌
技术层面：ARM工程师Jon Masters指出，ARM的Large System Extensions已经在硬件层面实现了远距离原子操作，说明业内早有类似探索。有人联想到IBM AS/400的全地址空间映射，也有人提到伯克利的IRAM项目（把处理器和内存做在一起）。还有人认为这就像Oracle数据库用的闪存加速器。
现实层面：集中在三个问题。第一是钱——有人猜这东西要2.5万美元，DDR5本身已经够贵了。第二是速度权衡，多一层内存就多一层延迟，不是所有场景都赚。第三是适用范围，有人认为这主要是超算用的，离普通电脑还远，乐观估计要4年。
对比与感慨：多人提到英特尔傲腾——如果傲腾还活着，今天可能就是它的高光时刻。也有人拿自己笔记本举例，说大SSD做swap已经很爽了，CXL是更极致的版本。
氛围：有人兴奋到“快拿走我的钱”，有人担心“内存脱离CPU指导”是不是有点吓人。