内存扩展多年来一直很受欢迎，尤其是因为机器学习模型对内存容量的需求不断增加。为此，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个独立的作业。内部的网络连接子系统到L3缓存和内存。两个Arm Cortex A53核心处理控制功能。MX1采用三星的4nm工艺制造，功耗为40W，这意味着每个RISC-V核心的功耗不到13mW。考虑到来自4个DIMM的功耗，板子的功耗为90W。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内存给予特殊处理，较大的页面大小对于巨大的内存块可能是有意义的。