GoKawiilMemory expansion has been attractive for years, and is more relevant than ever because ML models have an insatiable appetite for memory capacity. In response, XCENA has worked with Samsung to create a CXL memory expansion device that can also host SSDs and do compute. The device is called MX1, where MX stands for “Memory Xcelerator”.
On the memory expansion front, the MX1 can host up to 2 TB of DDR5 memory and connects to the host via a PCIe 6/CXL 3.2 x8 interface. The MX1 therefore has 128 GB/s of bandwidth to the host, or 64 GB/s in each direction. XCENA exposes another eight downstream PCIe 6 lanes that can be used to connect SSDs. SSD storage can be exposed as memory, with the MX1’s attached DRAM acting as a cache. The combination of DDR5 slots and downstream PCIe lanes lets the MX1 make a massive amount of memory capacity visible to the host. However, MX1’s most exciting feature is arguably a substantial amount of onboard compute.
MX1’s chip hosts 3072 RISC-V cores, divided into clusters of 32 cores that share L2 caches and a data TLB. Sets of four clusters form a “subsystem”, which acts as the smallest job allocation unit. MX1 has 24 subsystems, letting it run 24 independent jobs at a time. An in-house NoC connects the subsystems to L3 cache and memory. Two Arm Cortex A53 cores handle control functions. The MX1 is fabricated using Samsung’s 4nm process and consumes 40 W, implying that each each RISC-V core draws a bit under 13 mW. The board consumes 90W when accounting for power consumption from 4 DIMMs.
XCENA uses a lot of small cores because they’re targeting data-parallel workloads, where single threaded performance is less important than utilizing high memory bandwidth while maximizing power efficiency. It’s a strategy with parallels to Intel’s Xeon Phi, which similarly used a large number of low-clocked, relatively weak cores to take on highly parallel tasks.
Cache hierarchy within a cluster
Each RISC-V core uses in-order execution and runs at a pedestrian 1.1 GHz. XCENA’s cache hierarchy is almost GPU-like because it tries to avoid address translation overhead and varies cache sharing wildly at the upper levels. Each core has a 4 KB virtually addressed L1 data cache. Data-side memory accesses don’t go through address translation unless they miss L1D. 128 KB L2 data caches are shared across a cluster, as are TLBs for accelerating address translation. The L2 data cache is virtually indexed and physically tagged (VIPT), much like L1D caches in many conventional CPUs. From the instruction side, sets of four RISC-V cores share a 8 KB instruction cache. XCENA aims to contain kernel hot loops within this instruction cache, while a cluster-level 128 KB L2 instruction cache handles larger instruction footprints.
Instruction fetches operate directly on physical addresses and do not use virtual memory. Instruction accesses therefore don’t need address translation or TLBs. XCENA sets aside predefined device physical addresses for code, and restricts the program counter to those regions. Doing so prevents the RISC-V cores from accidentally jumping to data. XCENA handles process level isolation by isolating jobs at subsystem (128-core) boundaries. Presumably they also divide up the code region to give each subsystem its own code segment, which would prevent one process from accidentally executing another’s code.
From the MX1 product brief. MX1 is offered as a PCIe add-in card. The connector on the right connects to downstream PCIe SSDs
MX1’s programming model has parallels to OpenCL or CUDA. One kernel gets invoked many times, and each invocation uses an index to figure out what data it should process. Specifically, mu::getTaskIdx() is analogous to OpenCL’s get_global_id(). This model encourages code sharing across many cores, so sharing L1 instruction caches makes sense. If a loop is small enough, the four cores that share an instruction cache may fetch the same address at the same time, likely letting the instruction cache satisfy multiple fetches with a broadcast read.
On the data side, MX1 uses virtual memory and operates on the same virtual addresses as host code. Host and MX1 code can therefore share pointers, much like with OpenCL’s SVM. XCENA’s software sets up page tables to maintain the same mappings as the host’s. Each RISC-V core has a virtually addressed 4 KB L1 data cache, letting the core avoid address translation on a L1D hit. The cluster-level L2 cache is virtually addressed and physically tagged (VIPT), much like L1D caches on many CPUs. L2 indexing proceeds in parallel with a lookup in the cluster-shared TLB. The TLB has 1024 entries for 64 KB pages, and 8 entries for 1 GB pages. 64 KB pages allow more TLB coverage than typical 4 KB pages, but operating systems tend to use smaller page sizes to reduce overhead when paging to disk, copying pages, or clearing pages. However, XCENA expects the OS to give CXL memory special treatment, and larger page sizes may make sense for a giant block of expansion memory. 64 KB pages also align with typical SSD block sizes. As an aside, using different sets of TLB entries for different page sizes makes sense because it lets the TLB use different indexing schemes for each page size.
... continue reading