Hot Chips 2026：XCENA 与三星的近存计算 CXL 设备 MX1

CXL内存扩展，外加一点计算能力 Chester Lam 2026年8月30日

多年来，内存扩展一直颇具吸引力，如今更是比以往任何时候都更相关，因为机器学习模型对内存容量的需求永无止境。为此，XCENA与三星合作打造了一款CXL内存扩展设备，它还能承载SSD并执行计算。该设备名为MX1，其中MX代表“Memory Xcelerator”（内存加速器）。

在内存扩展方面，MX1最多可承载2 TB的DDR5内存，并通过PCIe 6/CXL 3.2 x8接口连接主机。因此，MX1到主机的带宽为128 GB/s，即每个方向64 GB/s。XCENA还额外引出八条下行PCIe 6通道，可用于连接SSD。SSD存储可以作为内存暴露出来，由MX1所连接的DRAM充当缓存。DDR5插槽与下行PCIe通道的结合，使MX1能够向主机呈现海量的内存容量。然而，MX1最令人兴奋的特性可以说是其可观的板载计算能力。

MX1的芯片集成3072个RISC-V核心，分为32核一组的集群，集群内共享L2缓存和数据TLB。每四个集群组成一个“子系统”，子系统是最小的任务分配单元。MX1共有24个子系统，可同时运行24个独立任务。自研的片上网络（NoC）将各子系统连接到L3缓存和内存。两个Arm Cortex A53核心负责控制功能。MX1采用三星4nm工艺制造，功耗为40 W，这意味着每个RISC-V核心的功耗略低于13 mW。计入4条DIMM的功耗后，整板功耗为90 W。

XCENA使用大量小核心，是因为其面向的是数据并行工作负载，在这类负载中，单线程性能不如在最大化能效的同时利用高内存带宽那么重要。这一策略与英特尔的至强融核（Xeon Phi）有相似之处，后者同样使用大量低频率、相对较弱的核心来处理高度并行的任务。

集群内的缓存层级

每个RISC-V核心采用顺序执行，运行频率仅为1.1 GHz。XCENA的缓存层级几乎类似GPU，因为它试图避免地址转换开销，并且在上层缓存中共享方式变化很大。每个核心都有一个4 KB的虚拟寻址L1数据缓存。数据侧的内存访问只有在L1D未命中时才会经过地址转换。128 KB的L2数据缓存在一个集群内共享，用于加速地址转换的TLB也是如此。L2数据缓存采用虚拟索引、物理标记（VIPT），与许多传统CPU中的L1D缓存非常相似。在指令侧，每四个RISC-V核心共享一个8 KB指令缓存。XCENA的目标是将内核热循环容纳在该指令缓存内，而集群级的128 KB L2指令缓存则负责处理更大的指令足迹。

指令取指直接基于物理地址进行，不使用虚拟内存。因此，指令访问不需要地址转换或TLB。XCENA为代码预留了预定义的设备物理地址，并将程序计数器限制在这些区域内。这样做可以防止RISC-V核心意外跳转到数据区。XCENA通过在子系统（128核）边界隔离任务来实现进程级隔离。据推测，他们还会划分代码区域，为每个子系统提供各自的代码段，从而防止一个进程意外执行另一个进程的代码。

摘自MX1产品简介。MX1以PCIe附加卡的形式提供。右侧的连接器用于连接下行PCIe SSD

MX1的编程模型与OpenCL或CUDA有相似之处。一个内核会被多次调用，每次调用都使用索引来确定应处理哪些数据。具体来说，mu::getTaskIdx()类似于OpenCL的get_global_id()。这种模型鼓励在许多核心之间共享代码，因此共享L1指令缓存是合理的。如果循环足够小，共享一个指令缓存的四个核心可能会同时取指同一地址，这很可能让指令缓存通过广播读取来满足多个取指请求。

在数据侧，MX1使用虚拟内存，并与主机代码操作相同的虚拟地址。因此，主机和MX1代码可以共享指针，这与OpenCL的SVM非常相似。XCENA的软件会设置页表，以维持与主机相同的映射。每个RISC-V核心都有一个虚拟寻址的4 KB L1数据缓存，使核心在L1D命中时无需进行地址转换。集群级L2缓存采用虚拟寻址、物理标记（VIPT），与许多CPU上的L1D缓存非常相似。L2索引与集群共享TLB中的查找并行进行。该TLB有1024个用于64 KB页面的条目，以及8个用于1 GB页面的条目。64 KB页面比典型的4 KB页面提供更大的TLB覆盖范围，但操作系统往往使用较小的页面大小，以减少换页到磁盘、复制页面或清除页面时的开销。不过，XCENA期望操作系统对CXL内存给予特殊处理，而对于一大块扩展内存来说，较大的页面大小可能是合理的。64 KB页面也与典型的SSD块大小一致。另外，为不同页面大小使用不同的TLB条目集合是合理的，因为这可以让TLB为每种页面大小使用不同的索引方案。

扩展RISC-V

XCENA利用RISC-V的可扩展性，在子系统层面实现了自定义的向量处理引擎（VPE）。每个RISC-V核心都有一个VPE命令队列，可以请求VPE加速各种向量操作。XCENA很可能使用特殊指令将消息入队到VPE命令队列，并期望代码将其视为一个巨大的共享协处理器。VPE支持FP32和FP16，在整个芯片上提供约3 TFLOPS的点积吞吐量。如果VPE与核心一样运行在1.1 GHz，那么每个VPE每周期可维持128 FLOPS。奇怪的是，VPE似乎并不加速整数运算。也许XCENA期望代码直接在RISC-V核心上运行整数运算。3072个RISC-V核心在1.1 GHz下，如果每个核心每周期完成一次运算，大约可获得每秒3T次整数运算。另一个奇怪之处是，XCENA的API通过返回错误码的内置函数来暴露VPU。代码必须显式检查溢出和无效访问等错误条件，这表明VPU指令不会引发异常。

SSD作为内存

MX1可以承载SSD，并将其作为CXL内存呈现给主机。XCENA称之为“无限内存”。MX1可以以RAID方式连接SSD，并且其下行PCIe和上行PCIe/CXL链路的带宽相匹配，这意味着它理论上可以仅靠SSD就饱和到主机的带宽。然而，与DRAM相比，SSD的延迟较高。MX1可以通过使用其连接的DDR5缓存SSD内容来缓解这种延迟。缓存以64 KB页面工作，并配有片上1024项映射缓存。映射缓存的作用类似于TLB，跟踪映射到SSD后端地址的DRAM页面。如果访问在映射缓存中未命中，就会引发缺页，由运行在MX1的RISC-V核心上的固件处理。固件通过从SSD获取数据并更新映射来处理缓存未命中。我感到困惑，因为1024项映射缓存以64 KB页面计算只能覆盖64 MB，但XCENA的文档表明缓存默认为16 GB，并且容量可按16 MB步进调整。鉴于映射缓存的结构，我不确定这是如何工作的。

为了在使用SSD作为内存时进一步利用连接的DDR5，用户可以配置“固定前缀”，将SSD后端地址固定到DRAM。固定前缀区域的大小可按16 MB步进配置。前缀意味着固定内存只能覆盖连续的地址空间，不具备页面级缓存的灵活性。因此，固定前缀内存最适合用于将频繁访问的缓冲区保留在DRAM中。XCENA的网站给出了一个示例：在231 GB连接的DRAM中，有115.5 GB的固定内存。

对于共享文档场景，文档前缀被固定到DRAM

为了缓解SSD延迟，MX1还可以设置为从SSD预取。预取有助于保持IO路径繁忙，并有助于将数据获取与计算执行重叠。XCENA希望利用固定和预取来缓解SSD相对于DRAM的低性能。MX1还可以以RAID方式运行SSD以提高带宽。结语

MX1和LPDDR5X-PIM都将计算放置在靠近内存的位置，从而可以利用无法通过主机接口访问的高内部内存带宽。MX1比LPDDR5X-PIM更具说服力，因为它可以像具有自身板载内存的加速器一样工作。软件不必面对与PIM模式切换相关的权衡和复杂性，并且使用MX1的计算不会阻塞其他线程的内存访问。主机线程和MX1的RISC-V核心理论上可以使用锁等标准多线程技术在同一缓冲区上工作，以确保顺序。

CXL 4.0规范摘录

与缓存打交道也应该比LPDDR5X-PIM更直接。Type 3 CXL设备（CXL.mem）可以使用监听来反向无效化主机缓存行，使设备无需将内存区域设为不可缓存即可使其结果可见。同样，CXL.mem设备可以跟踪来自主机的读所有权请求，并了解其自身的计算核心何时可以安全地修改数据。我不确定MX1是否为其RISC-V核心充分利用了这一能力，但三星表示该设备可以使用CXL内存语义在CPU和GPU之间共享其内存。希望这意味着该设备可以使用监听来与处理器内存子系统集成。

与LPDDR5X-PIM一样，MX1的吞吐量并不算高。作为参照，一块板载内存带宽相近的Nvidia GeForce GTX 1080拥有8.8 TFLOPS的FP32算力，而MX1仅能提供3 TFLOPS。但原始吞吐量并非重点。MX1提供的是一种缓解内存扩展部分弊端的方法。MX1侧的计算可以避开受限的主机接口，并避免穿越CXL链路带来的功耗开销。一个假设拥有8.8 TFLOPS算力的加速器，如果必须通过CXL为每个FLOP加载一个字节，其算力将被限制在仅64 GFLOPS。在同样场景下，MX1即使同样未命中缓存，也能超过200 GFLOPS。

MX1的近存计算方案展现出前景。CXL内存扩展器到其DRAM池的带宽往往高于CXL链路所能承载的带宽。常规主机访问会让其中很大一部分带宽闲置，如果主机无法从缓存中满足足够多的访问请求，就可能导致主机侧计算受限于内存带宽。延迟也是一个问题，因为扩展内存的访问延迟会高于直连内存。近存计算方案可以缓解这些内存扩展问题，同时避免像LPDDR5X-PIM那样在常规内存控制器之下尝试同样做法所带来的弊端。

我希望看到各公司未来继续探索内存扩展器和近存计算。也许在遥远的未来，这项技术能够下放到消费级产品。我们很多人手头都有以往装机留下的闲置DRAM，会非常有兴趣用它来抵消购买新内存的成本。额外附带一些算力也无妨。