삼성과 XCENA의 MX1: 메모리 옆에 연산을 심은 CXL 확장 장치

메모리 용량은 오랫동안 매력적인 확장 지점이었지만, 거대 ML 모델이 끝없이 메모리를 요구하면서 그 중요성은 어느 때보다 커졌다. 이번 Hot Chips 2026에서 공개된 XCENA와 삼성의 MX1은 이 흐름에 대한 하나의 답이다. MX1은 'Memory Xcelerator'의 약자로, CXL 기반 메모리 확장 장치이면서 동시에 SSD를 붙일 수 있고, 무엇보다 상당한 규모의 온보드 연산 능력을 품고 있다는 점이 특징이다. 단순히 용량을 늘리는 데 그치지 않고 데이터가 놓인 자리에서 직접 계산을 수행한다는 발상이다.

메모리 확장 측면에서 MX1은 최대 2TB의 DDR5를 탑재하며, 호스트와는 PCIe 6/CXL 3.2 x8 인터페이스로 연결된다. 이 링크의 대역폭은 양방향 합산 128GB/s, 방향당 64GB/s다. 여기에 별도의 다운스트림 PCIe 6 레인 8개를 노출해 SSD를 연결할 수 있고, 이 SSD를 메모리처럼 호스트에 보여주면서 장착된 DRAM을 캐시로 활용한다. XCENA는 이 구성을 'Infinite Memory'라고 부른다. 다운스트림과 업스트림 대역폭을 맞춰 두었기 때문에 이론적으로는 SSD만으로도 호스트 링크 대역폭을 채울 수 있다는 설명이다.

작은 코어를 대량으로 모은 이유

MX1의 핵심은 3072개에 달하는 RISC-V 코어다. 32개 코어가 L2 캐시와 데이터 TLB를 공유하는 클러스터를 이루고, 클러스터 4개가 모여 하나의 '서브시스템'이 된다. 서브시스템은 작업 할당의 최소 단위이며, 총 24개가 있어 동시에 24개의 독립 작업을 처리한다. 제어는 두 개의 Arm Cortex-A53 코어가 맡고, 자체 설계한 NoC가 서브시스템과 L3 캐시·메모리를 잇는다. 삼성 4nm 공정으로 만들어졌고 칩 소비 전력은 40W로, 코어 하나당 13mW를 밑도는 수준이다. DIMM 4개까지 포함한 보드 전체는 90W를 소비한다.

XCENA가 느리고 약한 코어를 대량으로 묶은 이유는 명확하다. 단일 스레드 성능보다 높은 메모리 대역폭 활용과 전력 효율이 중요한 데이터 병렬 워크로드를 겨냥했기 때문이다. 이는 다수의 저클럭 약소 코어로 병렬 작업을 노렸던 인텔 Xeon Phi와 닮은 전략이다. 각 코어는 인오더 실행에 1.1GHz로 동작한다. 캐시 계층은 다소 GPU에 가깝다. 코어마다 4KB의 가상 주소 L1 데이터 캐시를 두어 L1D 히트 시에는 주소 변환을 건너뛰고, 클러스터 단위로 128KB VIPT L2 데이터 캐시와 TLB를 공유한다. 명령어 측은 코어 4개가 8KB 명령어 캐시를 나눠 쓰며 커널의 핫 루프를 이 안에 담는 것을 목표로 한다.

명령어 인출은 물리 주소로만 이뤄져 TLB가 필요 없고, 프로그램 카운터는 미리 정해둔 코드 영역으로 제한된다. 코어가 실수로 데이터 영역으로 점프하는 것을 막기 위한 설계다. 프로세스 격리는 128코어 단위인 서브시스템 경계에서 처리된다. 프로그래밍 모델은 OpenCL이나 CUDA와 유사해, 하나의 커널이 여러 번 호출되고 각 호출이 인덱스로 자기가 처리할 데이터를 판별한다. mu::getTaskIdx()가 OpenCL의 get_global_id()에 대응한다. 데이터 측은 호스트와 동일한 가상 주소를 공유하므로, OpenCL의 SVM처럼 호스트와 MX1 코드가 포인터를 주고받을 수 있다.

벡터 엔진과 SSD 캐싱의 실제

연산 가속은 RISC-V의 확장성을 활용한 커스텀 벡터 처리 엔진(VPE)이 담당한다. 서브시스템 단위로 배치되며 각 코어가 명령 큐를 통해 벡터 연산을 요청하는, 일종의 공유 코프로세서 구조다. FP32와 FP16을 지원하고 칩 전체에서 약 3TFLOPS의 내적 처리량을 낸다. 흥미롭게도 정수 연산은 가속하지 않는데, 정수 작업은 RISC-V 코어에서 직접 돌리는 것을 전제한 것으로 보인다. 또한 VPE API는 오버플로나 잘못된 접근 같은 오류를 예외가 아니라 반환 코드로 알려, 코드가 명시적으로 오류를 확인해야 한다.

SSD를 메모리로 쓸 때의 지연은 DDR5 캐싱으로 완화한다. 캐싱은 64KB 페이지 단위로 이뤄지고 온칩 1024엔트리 맵 캐시가 TLB처럼 매핑을 추적한다. 미스가 나면 페이지 폴트가 발생하고 RISC-V 코어에서 도는 펌웨어가 SSD에서 데이터를 가져와 매핑을 갱신한다. 다만 1024엔트리로는 64KB 페이지 기준 64MB밖에 못 덮는데 문서상 기본 16GB에 16MB 단위 조정이 가능하다고 해, 구조상 앞뒤가 맞지 않는 대목도 있다. 자주 쓰는 버퍼는 'pinned prefix'로 DRAM에 고정할 수 있으며, XCENA는 231GB DRAM 중 115.5GB를 고정한 예시를 든다.

실무적으로 MX1을 GPU 같은 가속기로 착각하면 안 된다. 비슷한 온보드 대역폭을 가진 GTX 1080이 8.8TFLOPS의 FP32를 내는 데 비해 MX1은 3TFLOPS에 불과하다. 그러나 요점은 원시 처리량이 아니라 메모리 확장의 약점을 줄이는 데 있다. CXL 확장 메모리는 DRAM 풀 자체 대역폭이 호스트 링크보다 넓은 경우가 많아, 호스트가 접근하는 것만으로는 상당한 대역폭이 놀게 된다. FLOP당 1바이트를 CXL로 실어 나른다고 가정하면 8.8TFLOPS 가속기도 64GFLOPS로 묶이지만, 링크를 건너지 않는 MX1은 같은 조건에서 200GFLOPS를 넘길 수 있다. 여기에 CXL.mem의 스누프 기반 백무효화로 캐시 일관성을 다루기 쉽고, 삼성은 이 장치가 CPU와 GPU 사이에서 메모리를 공유할 수 있다고 밝힌다. LPDDR5X-PIM처럼 메모리 컨트롤러 아래에서 무리하게 연산을 넣지 않고, 별도의 온보드 메모리를 가진 가속기처럼 동작한다는 점이 MX1을 더 현실적인 선택지로 만든다.