MX1亮相，这款CXL设备实现近内存计算

内存扩展多年来一直备受关注，如今重要性更是达到新高度，因为机器学习模型对内存容量的需求几乎没有上限。为此，XCENA与三星合作推出一款CXL内存扩展设备，该设备还可挂载SSD并具备计算能力，产品命名为MX1，MX代表 “Memory Xcelerator（内存加速器）”。

在内存扩展方面，MX1最高可搭载2TB DDR5内存，通过PCIe 6/CXL 3.2 x8接口连接主机。MX1与主机之间带宽可达128GB/s，双向各64GB/s。XCENA还引出另外8路下行PCIe 6通道，可用于外接SSD。SSD存储可以被映射为内存，MX1板载的DRAM充当缓存。DDR5内存插槽加上下行PCIe通道，让MX1能够向主机暴露出超大内存容量。不过MX1最亮眼的特性，当属板载的大量计算算力。

MX1芯片集成3072个RISC‑V内核，每32个内核组成一个集群，共享L2缓存与数据TLB。4个集群构成一个 “子系统”，也是最小的任务分配单元。MX1一共拥有24个子系统，可同时运行24个独立任务。自研片上网络NoC将各个子系统连接至L3缓存与内存。两颗Arm Cortex‑A53内核负责控制管理功能。MX1采用三星4nm工艺制造，芯片功耗40W，折算下来每个RISC‑V内核功耗略低于13mW。算上4根DIMM内存的功耗，整板功耗为90W。

XCENA选用大量小内核，面向数据并行工作负载。这类场景下单线程性能并不重要，重点是充分利用高内存带宽，同时最大化能效。该思路与英特尔至强融核（Xeon Phi）类似，同样依靠大量低频、算力不算强的内核来处理高度并行任务。

每个RISC‑V内核采用顺序执行，运行频率仅1.1GHz。XCENA的缓存层级设计接近GPU，目的是降低地址转换开销，高层缓存采用差异化共享机制。每个内核配备4KB虚拟寻址L1数据缓存。只有L1数据缓存未命中时，数据访问才会触发地址转换。一个集群内共享128KB L2数据缓存以及用于加速地址转换的TLB。L2数据缓存采用虚拟索引、物理标记（VIPT），和很多传统CPU的L1数据缓存设计一致。指令侧，每4个RISC‑V内核共享8KB指令缓存。XCENA力求把内核热点循环放进该指令缓存；更大的指令占用，则由集群级128KB L2指令缓存承接。

指令读取直接使用物理地址，不使用虚拟内存。因此指令访问无需地址转换，也不依赖TLB。XCENA为代码预留设备专用物理地址区间，程序计数器被限定在这些地址范围内，避免RISC‑V内核意外跳转到数据地址。XCENA以子系统（128核）为边界实现任务之间的进程隔离。推测它还会对代码地址空间做划分，给每个子系统分配独立代码段，防止一个进程意外执行另一个进程的代码。

出自MX1产品简介。MX1为PCIe插卡形态，右侧连接器用于外接下行PCIe SSD

MX1的软件开发套件SDK设计思路类似OpenCL或CUDA。同一个内核函数会被大量调用，每次调用依靠索引确定待处理的数据。具体来说，mu::getTaskIdx()对应OpenCL的get_global_id()。这套编程模型鼓励大量内核复用同一份代码，因此共享L1指令缓存具备实际意义。如果循环代码足够小，共享指令缓存的4个内核可能同时访问同一地址，指令缓存可以通过广播读取一次性响应多次取指请求。

数据侧，MX1使用虚拟内存，与主机代码使用相同虚拟地址。主机与MX1代码可以直接共用指针，类似OpenCL的SVM共享虚拟内存。XCENA软件会构建页表，维持与主机完全一致的地址映射。每个RISC‑V内核拥有4KB虚拟寻址L1数据缓存，L1命中时无需地址转换。集群级L2缓存为虚拟索引物理标记（VIPT），和很多CPU的L1数据缓存相同。L2索引查找与集群共享TLB查询并行执行。TLB包含1024项64KB页表项，以及8项1GB大页表项。相比常用的4KB页，64KB页可以提升TLB覆盖范围；但操作系统一般倾向使用更小页面，减少磁盘换页、页面拷贝、页面清零带来的开销。不过XCENA认为操作系统应当对CXL内存做特殊适配，大容量扩展内存场景适合使用大页，同时64KB页也和SSD典型块大小对齐。另外，不同页尺寸使用独立TLB表项，便于TLB为不同页面大小使用不同索引算法。

XCENA利用RISC‑V可扩展特性，在子系统层面实现了自研向量处理引擎VPE。每个RISC‑V内核拥有VPE命令队列，可以调度VPE完成向量加速任务。XCENA大概率通过自定义指令向VPE命令队列下发任务，将其作为大型共享协处理器使用。VPE支持FP32、FP16精度，整芯片点运算总吞吐量约3 TFLOPS。如果VPE和内核同样运行在1.1GHz，每个VPE每个周期可完成128次浮点运算。有意思的是VPE并不负责整数运算，推测整数运算直接由RISC‑V内核完成。3072颗1.1GHz的RISC‑V内核，假设每个周期完成一次整数运算，整数算力约3 TOPS。另外一处特点：XCENA的API通过带错误码的内置函数调用VPU，代码需要显式检查溢出、非法访问等错误，说明VPU指令不会触发硬件异常。

MX1可以挂载SSD，并向主机呈现为CXL内存，XCENA将该功能命名为 “无限内存（Infinite Memory）”。MX1支持SSD RAID模式，下行PCIe链路与上行PCIe/CXL链路带宽匹配，理论上仅依靠SSD就可以跑满主机侧带宽。但SSD相比DRAM存在较高访问延迟。MX1利用板载DDR5对SSD数据做缓存，以此降低延迟。缓存以64KB页为粒度，搭配片上1024项映射缓存。该映射缓存功能类似TLB，记录DRAM页与SSD后端地址的映射关系。如果映射缓存未命中，会触发缺页异常，由MX1上RISC‑V内核运行固件处理。固件从SSD读取数据并更新映射表，完成缓存缺失处理。这里存在一个疑问：1024项映射缓存，按64KB页计算仅能覆盖64MB地址空间，但官方文档显示该缓存默认配置为16GB，可按16MB步长调整容量，以该映射缓存架构来看，其实现原理尚不明确。

为了在SSD作为内存时进一步发挥板载DDR5的能力，用户可以配置 “固定前缀（pinned prefix）”，把SSD后端对应的部分地址强制常驻DRAM。固定前缀区域大小以16MB为单位调整。该机制只能覆盖一段连续地址空间，不具备页级缓存的灵活度，适合把高频访问的缓冲区常驻在DRAM中。官方文档举例，231GB可用DRAM中，可以配置115.5GB作为固定常驻内存。

文档示例场景：文档前缀区域被固定常驻到DRAM

为缓解SSD延迟，MX1还支持SSD预取。预取可以保持IO通路忙碌，实现数据读取与计算执行时间重叠。XCENA希望依靠内存常驻与预取，弥补SSD相对DRAM的性能短板。MX1同时支持SSD RAID，提升带宽。

MX1与LPDDR5X‑PIM都采用近内存计算思路，充分利用主机接口无法触达的芯片内部高内存带宽。相比LPDDR5X‑PIM，MX1方案更具备现实说服力，它本身就是一块自带内存的加速器。软件无需面对PIM模式切换带来的取舍与复杂度，调用MX1算力也不会抢占其他线程的内存访问。理论上主机线程与MX1的RISC‑V内核可以操作同一块缓冲区，通过锁等标准多线程机制保证执行顺序。

CXL4.0规范片段

缓存交互也比LPDDR5X‑PIM更加简单。Type3 CXL设备（CXL.mem）支持探听机制，主动使无效主机缓存行，设备更新的数据对外可见，不需要把内存区域设置为不可缓存。CXL.mem设备还可以追踪主机的读‑获取所有权请求，判断自身计算内核何时可以安全修改数据。尚不清楚MX1的RISC‑V内核是否完整利用这套能力，但三星表示该设备可以借助CXL内存语义，在CPU与GPU之间共享内存。意味着设备可以通过缓存探听，和处理器内存子系统协同工作。

和LPDDR5X‑PIM一样，MX1的绝对算力并不高。做个参照：Nvidia GeForce GTX 1080板载内存带宽和MX1接近，FP32算力可达8.8 TFLOPS，而MX1仅有3 TFLOPS。但原始算力并不是它的核心价值。MX1的意义在于缓解内存扩展带来的各类弊端。在MX1本地完成计算，可以避开带宽受限的主机接口，省去跨CXL链路传输的功耗开销。假设一款8.8 TFLOPS的加速器，如果每一次浮点运算都要通过CXL读取1字节数据，实际算力会被限制到仅64 GFLOPS；而同样场景下，即便发生缓存未命中，MX1仍可跑出超过200 GFLOPS。

MX1这套近内存计算方案前景可观。CXL内存扩展器自身DRAM池的内部带宽，往往高于CXL链路对外带宽。普通主机访问只能用到其中一部分带宽；如果主机缓存无法命中，主机端计算就容易遭遇内存带宽瓶颈。扩展内存本身访问延迟也高于板载直连内存。近内存计算可以缓解上述内存扩展痛点，同时规避LPDDR5X‑PIM这类在传统内存控制器内部实现计算所带来的缺陷。

期待行业继续探索内存扩展器与近内存计算技术。或许在更远的未来，这类技术可以下沉到消费级市场。很多用户手里留有旧设备拆下的DRAM，如果可以利用它们减少新内存采购成本会非常有吸引力，如果还附带额外计算能力就更好了。