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

多年来，内存扩展一直很有吸引力，而且比以往任何时候都更加重要，因为机器学习模型对内存容量有着永不满足的需求。作为回应，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 公开了另外 8 个下游 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 的功耗，该主板的功耗为 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 对于 64 KB 页有 1024 个条目，对于 1 GB 页有 8 个条目。 64 KB 页面比典型的 4 KB 页面允许更多的 TLB 覆盖范围，但操作系统倾向于使用较小的页面大小来减少分页到磁盘、复制页面或清除页面时的开销。然而，XCENA 希望操作系统给予 CXL 内存特殊处理，更大的页面大小可能对巨大的 ex 块有意义。