Xcena和三星的近内存计算CXL设备

# Hot Chips 2026：XCENA与三星的近内存计算CXL设备 来源：https://chipsandcheese.com/p/hot-chips-2026-xcena-and-samsungs 内存扩展多年来一直备受关注，并且因为机器学习模型对内存容量有着永不满足的需求，如今比以往任何时候都更相关。为此，XCENA与三星合作开发了一款CXL内存扩展设备，该设备还能托管SSD并执行计算。该设备名为MX1，其中MX代表“内存加速器”。 [](https://substackcdn.com/image/fetch/$s_!rmOv!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2Fc4f3c070-692f-4557-9569-ea6f1d9fb689_1273x713.png) 在内存扩展方面，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最令人兴奋的功能或许是板载的大量计算能力。 [](https://substackcdn.com/image/fetch/$s_!MH3q!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F4ee7420b-9c57-4949-afe6-0c1a2e943dee_1273x714.png) MX1芯片承载3072个RISC-V核心，分为32个核心一组的集群，共享L2缓存和数据TLB。每四个集群构成一个“子系统”，作为最小的任务分配单元。MX1拥有24个子系统，可同时运行24个独立任务。一个内部片上网络将子系统连接到L3缓存和内存。两个Arm Cortex A53核心负责处理控制功能。MX1采用三星的4纳米工艺制造，功耗为40瓦，这意味着每个RISC-V核心的功耗略低于13毫瓦。当计算4条DIMM的功耗时，整块电路板的功耗为90瓦。 [](https://substackcdn.com/image/fetch/$s_!PVf6!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F598a6fbd-719c-4807-a833-bb723d1019f9_1142x632.png) XCENA使用大量小核心，因为他们的目标是数据并行工作负载，在这类负载中，单线程性能不如利用高内存带宽并最大化能效重要。这一策略与英特尔的Xeon Phi有相似之处，后者同样使用大量低时钟频率、相对较弱的核心来处理高度并行的任务。 [](https://substackcdn.com/image/fetch/$s_!Frc0!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F2f1b630a-bbba-4a07-b6df-30367f4b78cb_235x357.png) 集群内的缓存层次结构 每个RISC-V核心使用顺序执行，运行频率仅为1.1 GHz。XCENA的缓存层次结构几乎类似于GPU，因为它试图避免地址转换开销，并在较高层次大幅改变缓存共享方式。每个核心有一个4 KB的虚拟地址L1数据缓存。数据侧内存访问除非未命中L1D，否则不会经过地址转换。128 KB的L2数据缓存在一个集群内共享，用于加速地址转换的TLB也是共享的。L2数据缓存是虚拟索引物理标记的，与许多传统CPU中的L1D缓存非常相似。指令侧，每四个RISC-V核心共享一个8 KB的指令缓存。XCENA旨在将内核热循环包含在此指令缓存中，而集群级的128 KB L2指令缓存则处理更大的指令占用空间。 [](https://substackcdn.com/image/fetch/$s_!fugx!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F157dba42-7d8a-4ed5-a426-b80d48288673_1279x718.png) 指令获取直接在物理地址上操作，不使用虚拟内存。因此，指令访问不需要地址转换或TLB。XCENA为代码保留了预定义的设备物理地址，并将程序计数器限制在这些区域内。这样做可以防止RISC-V核心意外跳转到数据区域。XCENA通过在子系统边界（128核）隔离任务来处理进程级隔离。他们可能还划分了代码区域，为每个子系统分配自己的代码段，这可以防止一个进程意外执行另一个进程的代码。 [](https://substackcdn.com/image/fetch/$s_!PIRm!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F65c47d03-24ab-459b-925b-6ae26abccce6_663x289.png) 摘自MX1产品简报。MX1作为PCIe附加卡提供。右侧连接器连接到下行PCIe SSD。 MX1的编程模型（https://xcena.com/sdk）与OpenCL或CUDA有相似之处。一个内核被多次调用，每次调用使用一个索引来确定它应该处理哪些数据。具体来说，mu::getTaskIdx()类似于OpenCL的get_global_id()。这种模型鼓励跨多个核心的代码共享，因此共享L1指令缓存是有意义的。如果一个循环足够小，共享指令缓存的四个核心可能同时获取相同的地址，这很可能使指令缓存能够通过广播读取满足多次获取。 [](https://substackcdn.com/image/fetch/$s_!2XTO!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F6fda6a6c-a5d5-4b91-af10-d308b17f8bcf_1274x717.png) 在数据侧，MX1使用虚拟内存，并与主机代码使用相同的虚拟地址。因此，主机和MX1代码可以共享指针，类似于OpenCL的SVM。XCENA的软件设置页表以保持与主机相同的映射。每个RISC-V核心有一个虚拟地址的4 KB L1数据缓存，使核心在L1D命中时可以避免地址转换。集群级L2缓存是虚拟寻址物理标记的，与许多CPU上的L1D缓存非常相似。L2索引与集群共享TLB的查找并行进行。TLB有1024个条目用于64 KB页面，8个条目用于1 GB页面。64 KB页面比典型的4 KB页面提供更多的TLB覆盖范围，但操作系统倾向于使用较小的页面大小，以减少换出到磁盘、复制页面或清空页面时的开销。然而，XCENA期望操作系统对CXL内存给予特殊处理，对于巨大的扩展内存块，较大的页面大小可能是合理的。64 KB页面也与典型的SSD块大小对齐。顺便提一下，为不同页面大小使用不同的TLB条目组是有意义的，因为它允许TLB为每种页面大小使用不同的索引方案。 [](https://substackcdn.com/image/fetch/$s_!S5gM!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F14fe4378-3157-4751-892b-030d082a1032_1276x714.png) XCENA利用RISC-V的可扩展性在子系统级别实现了一个定制的向量处理引擎。每个RISC-V核心有一个VPE命令队列，可以要求VPE加速各种向量操作（https://xcena-dev.github.io/mu_lib/cpp/classmu_1_1vector_1_1_vdma_error.html#details）。很可能，XCENA使用特殊指令将消息排入VPE命令队列，并期望代码将其视为一个巨大的共享协处理器。VPE支持FP32和FP16，并在整个芯片上提供约3 TFLOPS的点积吞吐量。如果VPE与核心一样运行在1.1 GHz，那么每个VPE每个周期可以持续执行128次浮点运算。奇怪的是，VPE似乎不加速整数操作。也许XCENA期望代码直接在RISC-V核心上运行整数操作。3072个RISC-V核心在1.1 GHz下，如果每个核心每个周期完成一个操作，将获得大约3万亿次整数操作/秒。另一个奇特之处是，XCENA的API通过返回错误代码的内置函数暴露VPU。代码必须显式检查溢出和无效访问等错误条件，这表明VPU指令不会引发异常。 MX1可以托管SSD，这些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的文档（https://xcena-dev.github.io/InfiniteMemory_docs/cli.html）表明缓存默认为16 GB，并且容量可以以16 MB为步长调整。我不确定在给定映射缓存结构的情况下这是如何工作的。 [](https://substackcdn.com/image/fetch/$s_!c_0T!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2Fd7646c81-afbd-4657-8476-d7a876bd1d67_1277x718.png) 为了进一步利用使用SSD作为内存时的附带DDR5，用户可以配置“固定前缀”，将SSD支持的地址固定在DRAM中。固定前缀区域大小可以按16 MB步长配置。前缀意味着固定内存只能覆盖一个连续的地址空间，不具备页面级缓存的灵活性。因此，固定前缀内存最适合用于将频繁访问的缓冲区保留在DRAM中。XCENA的网站（https://xcena-dev.github.io/InfiniteMemory_docs/cli.html）给出了一个示例，在231 GB的附带DRAM中有115.5 GB的固定内存。 [](https://substackcdn.com/image/fetch/$s_!UquA!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2Fec9cc8ff-87eb-415a-bb8b-f4b42baa4567_1279x718.png) 对于共享文档的情况，文档前缀被固定在DRAM中。 为了缓解SSD延迟，MX1还可以设置为从SSD预取。预取有助于保持I/O路径繁忙，并有助于将数据获取与计算执行重叠。XCENA希望通过固定和预取来缓解SSD与DRAM相比性能较低的问题。MX1还可以以RAID方式运行SSD以增加带宽。 MX1和LPDDR5X-PIM都将计算放置在内存附近，它们可以利用无法通过主机接口访问的高内部内存带宽。MX1比LPDDR5X-PIM提出了一个更令人信服的案例，因为它可以充当具有自己板载内存的加速器。软件不必面对与PIM模式切换相关的权衡和复杂性，并且使用MX1的计算不会阻止来自其他线程的内存访问。主机线程和MX1 RISC-V核心理论上可以处理相同的缓冲区，使用标准的如锁之类的多线程技术来确保顺序。 [](https://substackcdn.com/image/fetch/$s_!zCTg!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2Fff76197c-8ca1-4828-9061-b28502415c95_549x624.png) 摘自CXL 4.0规范 与LPDDR5X-PIM相比，处理缓存也应该更直接。类型3 CXL设备（CXL.mem）可以使用窥探来使主机缓存行无效，让设备在不使内存区域不可缓存的情况下使其结果可见。类似地，CXL.mem设备可以跟踪来自主机的读取以获取所有权请求，并理解其自身的计算核心何时可以安全地修改数据。我不确定MX1是否完全利用了其RISC-V核心的这一功能，但三星表示该设备可以使用CXL内存语义在CPU和GPU之间共享内存。希望这意味着该设备可以使用窥探来与处理器内存子系统集成。 [](https://substackcdn.com/image/fetch/$s_!ddiD!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F41bb30eb-05c0-4c6d-8a8d-bc707260e2f4_1275x718.png) 与LPDDR5X-PIM一样，MX1没有提供很高的吞吐量。作为参考，具有相似板载内存带宽的Nvidia GeForce GTX 1080具有8.8 TFLOPS的FP32计算能力，而MX1提供的是3 TFLOPS。但原始吞吐量不是重点。相反，MX1提供了一种方法来缓解一些内存扩展的缺点。MX1侧的计算可以避免受限的主机接口，并避免遍历CXL链路的功耗开销。一个具有8.8 TFLOPS计算能力的假设加速器，如果必须通过CXL每浮点运算加载一个字节，则其性能将仅限于64 GFLOPS。MX1在相同场景下即使缓存未命中也可能超过200 GFLOPS。 [](https://substackcdn.com/image/fetch/$s_!YpuW!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2Fa90ec7fe-f1cb-401e-a8cd-fdd530a4696b_1277x716.png) MX1的近内存计算方法显示出前景。CXL内存扩展器对其DRAM池的带宽通常高于CXL链路所能处理的。常规主机访问会留下很大一部分带宽未使用，这可能导致主机端计算在主机无法从其缓存中服务足够访问时受到内存带宽限制。延迟也是一个问题，因为扩展内存的访问延迟高于直接连接的内存。近内存计算选项可以缓解这些内存扩展问题，而不会像LPDDR5X-PIM那样在正常内存控制器下尝试相同操作所带来的缺点。 [](https://substackcdn.com/image/fetch/$s_!xVci!,f_auto,q_auto:good,fl_progressive:steep/https%3A%2F%2Fsubstack-post-media.s3.amazonaws.com%2Fpublic%2Fimages%2F1d391b68-cdf1-4dae-bd9f-363df87a086a_1271x715.png) 我希望看到公司未来继续探索内存扩展器和近内存计算。也许在遥远的未来，这项技术可以普及到消费者。我们许多人都有来自之前装机的闲置DRAM，并且会非常有兴趣使用它们来缓解购买新内存的成本。额外的计算能力也不会有什么坏处。