XCENAとSamsungのニアメモリコンピューティングCXLデバイス

Hot Chips 2026: XCENAとSamsungのニアメモリコンピューティングCXLデバイス メモリ拡張にコンピューティング機能を添えて Chester Lam 2026年8月30日 82シェア メモリ拡張は何年も前から魅力的でしたが、MLモデルがメモリ容量を飽くほど求めるようになったため、今やこれまで以上に重要になっています。これに対応するため、XCENAはSamsungと協力して、SSDとコンピューティングもホストできるCXLメモリ拡張デバイスを作成しました。このデバイスはMX1と呼ばれ、MXは「Memory Xcelerator」の略です。メモリ拡張の面では、MX1は最大2TBのDDR5メモリを搭載でき、PCIe 6/CXL 3.2 x8インターフェースを介してホストに接続されます。したがって、MX1はホストへの帯域幅が128 GB/s、または双方向で64 GB/sです。XCENAは、SSDを接続するために使用できる8つの下流PCIe 6レーンをさらに公開しています。SSDストレージはメモリとして公開でき、MX1に接続されたDRAMはキャッシュとして機能します。DDR5スロットと下流PCIeレーンの組み合わせにより、MX1はホストに大量のメモリ容量を可視化できます。しかし、MX1の最もエキサイティングな機能は、おそらくかなりの量のオンボードコンピューティングです。MX1のチップは3072個のRISC-Vコアを搭載しており、L2キャッシュとデータTLBを共有する32コアのクラスターに分割されています。4つのクラスターのセットが「サブシステム」を形成し、これが最小のジョブ割り当て単位として機能します。MX1には24のサブシステムがあり、一度に24個の独立したジョブを実行できます。社内製NoCがサブシステムをL3キャッシュとメモリに接続します。2つのArm Cortex A53コアが制御機能を担当します。MX1はSamsungの4nmプロセスで製造され、40Wを消費します。これは、各RISC-Vコアが13mW未満を消費することを意味します。ボードは、4つのDIMMからの消費電力を考慮すると90Wを消費します。XCENAは、単一スレッドのパフォーマンスよりも高いメモリ帯域幅を利用し、電力効率を最大化することが重要なデータ並列ワークロードをターゲットにしているため、多くの小さなコアを使用しています。これは、IntelのXeon Phiの戦略に似ており、同様に多数の低クロックで比較的弱いコアを使用して、高度に並列化されたタスクに対応していました。クラスター内のキャッシュ階層各RISC-Vコアはインオーダー実行を使用し、1.1 GHzという控えめなクロックで動作します。XCENAのキャッシュ階層は、アドレス変換オーバーヘッドを回避しようとし、上位レベルでキャッシュ共有を大幅に変更するため、ほぼGPUに似ています。各コアには4KBの仮想アドレス指定可能なL1データキャッシュがあります。データサイドのメモリアクセスは、L1Dをヒットしない限り、アドレス変換を経由しません。128KBのL2データキャッシュはクラスター全体で共有され、アドレス変換を高速化するためのTLBも同様です。L2データキャッシュは、多くの従来のCPUのL1Dキャッシュと同様に、仮想インデックス付けされ、物理タグ付けされています（VIPT）。命令サイドでは、4つのRISC-Vコアのセットが8KBの命令キャッシュを共有します。XCENAは、カーネルのホットループをこの命令キャッシュ内に収めることを目指しており、クラスターレベルの128KB L2命令キャッシュは、より大きな命令フットプリントに対応します。命令フェッチは物理アドレスに直接作用し、仮想メモリを使用しません。したがって、命令アクセスではアドレス変換やTLBは必要ありません。XCENAは、コードのために定義済みのデバイス物理アドレスを確保し、プログラムカウンタをそれらの領域に制限します。これにより、RISC-Vコアが誤ってデータにジャンプするのを防ぎます。XCENAは、サブシステム（128コア）の境界でジョブを分離することにより、プロセスレベルの分離を処理します。おそらく、彼らはコード領域も分割して、各サブシステムに独自のコードセグメントを提供し、あるプロセスが別のプロセスのコードを誤って実行するのを防ぐでしょう。MX1製品ブリーフより。MX1はPCIeアドインカードとして提供されます。右側のコネクタは、下流PCIe SSDに接続されます。MX1のプログラミングモデルはOpenCLまたはCUDAに似ています。1つのカーネルが何度も呼び出され、各呼び出しはインデックスを使用して処理すべきデータを特定します。具体的には、mu::getTaskIdx()はOpenCLのget_global_id()に相当します。このモデルは、多数のコア間でのコード共有を促進するため、L1命令キャッシュの共有は理にかなっています。ループが十分に小さい場合、命令キャッシュを共有する4つのコアは同時に同じアドレスをフェッチする可能性があり、命令キャッシュはブロードキャスト読み取りで複数のフェッチを満足させることができます。データサイドでは、MX1は仮想メモリを使用し、ホストコードと同じ仮想アドレスで動作します。したがって、ホストとMX1のコードは、OpenCLのSVMと同様にポインタを共有できます。XCENAのソフトウェアは、ホストと同じマッピングを維持するためにページテーブルを設定します。各RISC-Vコアには仮想アドレス指定可能な4KB L1データキャッシュがあり、コアはL1Dヒット時にアドレス変換を回避できます。クラスターレベルのL2キャッシュは仮想アドレス指定され、物理タグ付けされています（VIPT）。これは、多くのCPUのL1Dキャッシュに似ています。L2インデックス付けは、クラスター共有TLBのルックアップと並行して進行します。TLBは、64KBページ用に1024エントリ、1GBページ用に8エントリを持っています。64KBページは、通常の4KBページよりも多くのTLBカバレッジを提供しますが、オペレーティングシステムは、ディスクへのページング、ページのコピー、またはページのクリア時のオーバーヘッドを削減するために、より小さいページサイズを使用する傾向があります。しかし、XCENAはOSがCXLメモリに特別な扱いをすることを期待しており、より大きなページサイズは巨大な拡張メモリブロックに適している可能性があります。64KBページは、通常のSSDブロックサイズとも一致します。余談ですが、異なるページサイズに異なるTLBエントリセットを使用することは、TLBが各ページサイズに異なるインデックス付けスキームを使用できるため、理にかなっています。RISC-Vの拡張XCENAは、RISC-Vの拡張性を利用して、サブシステムレベルでカスタムベクトル処理エンジン（VPE）を実装しています。各RISC-VコアはVPEコマンドキューを取得し、VPEにさまざまなベクトル操作を高速化するように要求できます。おそらく、XCENAは特別な命令を使用してVPEコマンドキューにメッセージをエンキューし、コードがそれを巨大な共有コプロセッサとして扱うことを期待しています。VPEはFP32とFP16をサポートし、チップ全体で約3 TFLOPSのドット積スループットを提供します。VPEがコアと同じ1.1 GHzで動作する場合、各VPEはサイクルあたり128 FLOPSを維持できます。興味深いことに、VPEは整数演算を高速化しないようです。おそらく、XCENAはコードが整数演算をRISC-Vコアで直接実行することを期待しています。1.1 GHzで3072個のRISC-Vコアは、各コアがサイクルあたり1つの演算を完了すると仮定すると、約3兆の整数演算/秒になります。もう一つの興味深い点は、XCENAのAPIが組み込み関数を使用してVPUを公開しており、エラーコードを返すことです。コードはオーバーフローや無効なアクセスなどのエラー条件を明示的にチェックする必要があり、VPU命令は例外を発生させないことを示唆しています。SSDをメモリとしてMX1はSSDをホストでき、これらはCXLメモリとしてホストに提示されます。XCENAはこれを「Infinite Memory」と呼んでいます。MX1はRAID構成でSSDを搭載でき、下流PCIeと上流PCIe/CXLリンクの両方に一致する帯域幅を持っているため、理論的にはSSDのみを使用してホストへの帯域幅を飽和させることができます。しかし、SSDはDRAMと比較してレイテンシが高いです。MX1は、接続されたDDR5を使用してSSDコンテンツをキャッシュすることで、そのレイテンシを軽減できます。キャッシュは64KBページで機能し、オンチップの1024エントリマップキャッシュを備えています。マップキャッシュはTLBのように機能し、SSDバックアドレスにマッピングされたDRAMページを追跡します。アクセスがマップキャッシュでミスした場合、ページフォルトが発生し、MX1のRISC-Vコアで実行されるファームウェアによって処理されます。ファームウェアは、SSDからデータをフェッチし、マッピングを更新することでキャッシュミスを処理します。1024エントリのマップキャッシュは64KBページで64MBしかカバーしないため、混乱しています。しかし、XCENAのドキュメントでは、キャッシュはデフォルトで16GBであり、16MBステップで容量を調整できると示唆しています。マップキャッシュの構造を考えると、それがどのように機能するのかは不明です。SSDをメモリとして使用する際に接続されたDDR5をさらに活用するために、ユーザーは「ピン留めされたプレフィックス」を設定できます。このプレフィックスでは、SSDバックアドレスがDRAMにピン留めされます。ピン留めされたプレフィックス領域のサイズは、16MBステップで設定できます。プレフィックスということは、ピン留めされたメモリは連続したアドレス空間しかカバーできず、ページレベルキャッシュの柔軟性がないことを意味します。したがって、ピン留めされたプレフィックスメモリは、頻繁にアクセスされるバッファをDRAMに保持するのに最適です。XCENAのサイトでは、例としてw