XCENA MX1: Bước đột phá trong công nghệ mở rộng bộ nhớ CXL với khả năng tính toán tích hợp

Trong bối cảnh các mô hình máy học (ML) ngày càng "khát" dung lượng bộ nhớ, việc mở rộng bộ nhớ đã trở thành một nhu cầu cấp thiết. Tại hội nghị Hot Chips 2026, XCENA hợp tác với Samsung đã giới thiệu MX1, một thiết bị mở rộng bộ nhớ CXL không chỉ đơn thuần cung cấp thêm RAM mà còn tích hợp một cụm tính toán mạnh mẽ với 3072 lõi RISC-V. Đây được xem là một giải pháp độc đáo cho bài toán tính toán gần bộ nhớ (near-memory compute), mở ra hướng đi mới cho các trung tâm dữ liệu.

Sơ đồ kiến trúc tổng quan của thiết bị MX1

Kiến trúc phần cứng và khả năng kết nối linh hoạt

Về mặt mở rộng bộ nhớ, MX1 có thể hỗ trợ tối đa 2 TB bộ nhớ DDR5 và kết nối với máy chủ thông qua giao diện PCIe 6/CXL 3.2 x8, cung cấp băng thông lên tới 128 GB/s (64 GB/s mỗi chiều). Điểm đặc biệt là MX1 còn cung cấp thêm 8 làn PCIe 6 downstream, cho phép kết nối các ổ SSD. Dung lượng từ SSD có thể được "phơi bày" như bộ nhớ thông thường, với DRAM gắn liền hoạt động như một lớp đệm (cache) tốc độ cao. Sự kết hợp giữa khe cắm DDR5 và các làn PCIe downstream giúp MX1 tạo ra một vùng bộ nhớ khổng lồ cho hệ thống máy chủ.

Tuy nhiên, tính năng ấn tượng nhất của MX1 nằm ở sức mạnh tính toán đi kèm. Con chip bên trong MX1 sở hữu 3072 lõi RISC-V, được chia thành các cụm (cluster) 32 lõi dùng chung bộ nhớ đệm L2 và TLB dữ liệu. Bốn cụm tạo thành một "hệ thống con" (subsystem) gồm 128 lõi, hoạt động như một đơn vị phân bổ công việc nhỏ nhất. Toàn bộ MX1 có 24 hệ thống con, cho phép chạy đồng thời 24 tác vụ độc lập. Ngoài ra, hai lõi Arm Cortex A53 đảm nhiệm chức năng điều khiển.

MX1 được sản xuất trên tiến trình 4nm của Samsung với mức tiêu thụ điện năng chỉ 40W cho chip, trong khi toàn bộ bo mạch (bao gồm cả 4 thanh DIMM) tiêu thụ khoảng 90W. Mỗi lõi RISC-V chỉ tốn chưa đến 13mW điện năng, cho thấy hiệu quả năng lượng vượt trội.

Chiến lược "nhiều lõi nhỏ" cho tác vụ song song dữ liệu

XCENA lựa chọn sử dụng một số lượng lớn các lõi nhỏ vì họ hướng đến các tác vụ song song dữ liệu (data-parallel workloads), nơi mà hiệu năng đơn luồng không quan trọng bằng việc tận dụng băng thông bộ nhớ cao và tối ưu điện năng. Chiến lược này có nhiều điểm tương đồng với dòng sản phẩm Intel Xeon Phi trước đây, vốn cũng sử dụng một lượng lớn lõi có xung nhịp thấp để xử lý các tác vụ song song hóa cao.

Cấu trúc bộ nhớ đệm trong một cụm lõi RISC-V của MX1

Các lõi RISC-V sử dụng kiến trúc thực thi theo thứ tự (in-order) với xung nhịp khiêm tốn 1.1 GHz. Điểm thú vị là kiến trúc bộ nhớ đệm của XCENA khá giống GPU, khi nó cố gắng giảm thiểu chi phí dịch địa chỉ (address translation). Mỗi lõi có một bộ đệm dữ liệu L1 dung lượng 4 KB được định địa chỉ ảo (virtually addressed). Các truy cập bộ nhớ không cần đi qua dịch địa chỉ nếu trúng L1D. Bộ đệm L2 128 KB được chia sẻ trong mỗi cụm, cùng với TLB để tăng tốc dịch địa chỉ. Đáng chú ý, L2 dữ liệu sử dụng kỹ thuật VIPT, tương tự như L1D trên các CPU thông thường.

Về phía lệnh, bốn lõi RISC-V chia sẻ một bộ đệm lệnh L1 8 KB. XCENA đặt mục tiêu chứa các vòng lặp nóng của kernel trong bộ đệm này, trong khi bộ đệm lệnh L2 128 KB ở cấp cụm sẽ xử lý các khối lệnh lớn hơn. Điểm đặc biệt là việc nạp lệnh hoạt động trực tiếp trên địa chỉ vật lý và không sử dụng bộ nhớ ảo, do đó không cần TLB. XCENA dành riêng một vùng địa chỉ vật lý thiết bị cho mã lệnh và giới hạn con trỏ lệnh trong vùng đó, ngăn các lõi RISC-V vô tình nhảy vào vùng dữ liệu. Sự cô lập ở mức tiến trình được thực hiện bằng cách cô lập công việc theo ranh giới của các hệ thống con 128 lõi.

Mô hình lập trình giống CUDA/OpenCL

Mô hình lập trình của MX1 có nhiều nét tương đồng với OpenCL hoặc CUDA. Một kernel được gọi nhiều lần, và mỗi lần gọi sử dụng một chỉ số để xác định dữ liệu cần xử lý. Hàm mu::getTaskIdx() của XCENA tương tự như get_global_id() của OpenCL. Mô hình này khuyến khích việc chia sẻ mã giữa nhiều lõi, do đó việc dùng chung bộ đệm lệnh L1 là hợp lý. Nếu một vòng lặp đủ nhỏ, bốn lõi dùng chung bộ đệm lệnh có thể cùng nạp một địa chỉ, cho phép bộ đệm phục vụ nhiều truy cập bằng một lần đọc broadcast.

Về phía dữ liệu, MX1 sử dụng bộ nhớ ảo và hoạt động trên cùng địa chỉ ảo với mã máy chủ. Nhờ đó, mã của máy chủ và mã của MX1 có thể dùng chung con trỏ, giống như cơ chế Shared Virtual Memory (SVM) của OpenCL. Mỗi lõi RISC-V có bộ đệm L1 4 KB định địa chỉ ảo, giúp tránh chi phí dịch địa chỉ khi trúng L1D. TLB dùng chung cấp cụm có 1024 mục cho trang 64 KB và 8 mục cho trang 1 GB. Việc sử dụng trang 64 KB cho phép vùng phủ TLB rộng hơn so với trang 4 KB thông thường, và kích thước này cũng khớp với kích thước block điển hình của SSD.

Mở rộng kiến trúc RISC-V với Vector Processing Engine

XCENA tận dụng khả năng mở rộng của kiến trúc RISC-V để triển khai một Vector Processing Engine (VPE) tùy chỉnh ở cấp hệ thống con. Mỗi lõi RISC-V có một hàng đợi lệnh VPE và có thể yêu cầu VPE tăng tốc các phép toán vector đa dạng. VPE hỗ trợ FP32 và FP16, cung cấp thông lượng khoảng 3 TFLOPS cho phép toán dot product trên toàn chip. Nếu VPE chạy ở 1.1 GHz, mỗi VPE có thể duy trì 128 FLOPS mỗi chu kỳ.

Điều thú vị là VPE dường như không tăng tốc các phép toán số nguyên. Có lẽ XCENA kỳ vọng các phép toán số nguyên sẽ chạy trực tiếp trên các lõi RISC-V. Với 3072 lõi ở 1.1 GHz, MX1 có thể đạt khoảng 3 nghìn tỷ phép toán số nguyên mỗi giây nếu mỗi lõi hoàn thành một phép toán mỗi chu kỳ. Một chi tiết đáng chú ý khác là API của XCENA hiển thị VPU thông qua các hàm tích hợp trả về mã lỗi, và mã phải kiểm tra tường minh các điều kiện lỗi như tràn số hoặc truy cập không hợp lệ, cho thấy lệnh VPU không tạo ra ngoại lệ.

Biến SSD thành bộ nhớ với "Infinite Memory"

Mô hình sử dụng bộ nhớ được ghim (pinned memory) cho tài liệu dùng chung

MX1 có khả năng chứa SSD và trình bày chúng cho máy chủ như bộ nhớ CXL, tính năng này được gọi là "Infinite Memory". Tuy nhiên, SSD có độ trễ cao hơn nhiều so với DRAM. Để giảm thiểu độ trễ này, MX1 sử dụng DDR5 gắn liền làm bộ nhớ đệm cho dữ liệu SSD. Cơ chế caching hoạt động với trang 64 KB và một bộ nhớ đệm ánh xạ (map cache) 1024 mục trên chip. Bộ đệm này hoạt động giống như TLB, theo dõi các trang DRAM được ánh xạ tới địa chỉ trên SSD. Nếu truy cập bị trượt trong map cache, một lỗi trang sẽ được xử lý bởi firmware chạy trên các lõi RISC-V của MX1.

Ngoài ra, MX1 còn cho phép cấu hình "pinned prefix" (vùng bộ nhớ ghim), nơi các địa chỉ SSD được ghim cố định vào DRAM. Kích thước vùng này có thể được cấu hình theo các bước 16 MB. Tính năng này phù hợp nhất để giữ một vùng đệm truy cập thường xuyên trong DRAM, như ví dụ với 115.5 GB bộ nhớ ghim trong tổng số 231 GB DRAM của XCENA. MX1 cũng hỗ trợ cơ chế prefetch từ SSD để giảm độ trễ, giúp giữ cho đường I/O luôn bận rộn và chồng lấp việc nạp dữ liệu với quá trình tính toán.

Nhận định và so sánh với các giải pháp khác

Cả MX1 và LPDDR5X-PIM của Samsung đều đặt khả năng tính toán gần bộ nhớ để khai thác băng thông nội bộ cao mà không thể truy cập qua giao diện máy chủ. Tuy nhiên, MX1 cho thấy một tiềm năng thuyết phục hơn vì nó có thể hoạt động như một bộ tăng tốc với bộ nhớ riêng. Phần mềm không phải đối mặt với sự đánh đổi và phức tạp liên quan đến chuyển đổi chế độ PIM, và việc sử dụng tính toán của MX1 sẽ không chặn các truy cập bộ nhớ từ các luồng khác.

Với các thiết bị Type 3 CXL (CXL.mem), MX1 có thể sử dụng cơ chế snoop để vô hiệu hóa các dòng cache của máy chủ, giúp kết quả tính toán của thiết bị được hiển thị mà không cần làm cho vùng bộ nhớ không thể cache. Samsung cũng xác nhận rằng thiết bị có thể sử dụng ngữ nghĩa bộ nhớ CXL để chia sẻ bộ nhớ của nó cho cả CPU và GPU.

Ví dụ về ngữ nghĩa bộ nhớ CXL trong truy cập chia sẻ

Về hiệu năng tuyệt đối, MX1 không mang lại thông lượng lớn. Để so sánh, một GPU Nvidia GeForce GTX 1080 với băng thông bộ nhớ tương đương có thể đạt 8.8 TFLOPS FP32, trong khi MX1 chỉ cung cấp 3 TFLOPS. Tuy nhiên, mục đích của MX1 không phải là hiệu năng thô. Thay vào đó, MX1 đưa ra một cách để giảm thiểu nhược điểm của việc mở rộng bộ nhớ. Việc tính toán trên MX1 giúp tránh giới hạn băng thông của giao diện CXL và tiết kiệm năng lượng vốn phải tiêu tốn khi truyền qua liên kết CXL. Trong kịch bản một bộ tăng tốc có 8.8 TFLOPS nhưng phải nạp 1 byte mỗi FLOP qua CXL sẽ chỉ đạt 64 GFLOPS, thì MX1 có thể vượt qua 200 GFLOPS trong cùng kịch bản ngay cả khi bị trượt cache.

Kết luận

Cách tiếp cận tính toán gần bộ nhớ của MX1 cho thấy nhiều hứa hẹn. Các bộ mở rộng CXL thường có băng thông đến vùng DRAM cao hơn những gì liên kết CXL có thể xử lý. Các truy cập thông thường của máy chủ sẽ bỏ phí một phần lớn băng thông đó, dẫn đến tình trạng máy chủ bị giới hạn băng thông bộ nhớ nếu không thể phục vụ đủ truy cập từ cache. Độ trễ cũng là một vấn đề vì bộ nhớ mở rộng có độ trễ cao hơn bộ nhớ gắn trực tiếp. Một tùy chọn tính toán gần bộ nhớ có thể giảm thiểu các vấn đề này mà không gặp phải nhược điểm của việc thực hiện điều tương tự dưới bộ điều khiển bộ nhớ thông thường như LPDDR5X-PIM.

Với sự phát triển không ngừng của các mô hình AI/ML, việc sở hữu những giải pháp mở rộng bộ nhớ thông minh và hiệu quả như MX1 sẽ ngày càng trở nên quan trọng. Có thể trong tương lai xa, công nghệ này sẽ được ứng dụng xuống phân khúc người dùng phổ thông, giúp tận dụng những thanh DRAM cũ từ các bản nâng cấp máy tính trước đây để giảm chi phí mua bộ nhớ mới. Một chút sức mạnh tính toán đi kèm cũng là một lợi ích không hề nhỏ.