Memory Architecture / Analysis

Apple のユニファイドメモリは
何を実現したのか

hUMA との比較による実態整理 ── 共有されているのは物理メモリであって、アドレス空間ではない

2026-08-20
対象:Apple Silicon (M1–M5) / Metal 3–4 / AMD hUMA・HSA
要旨

Apple はユニファイドメモリアーキテクチャ(UMA)を「CPU と GPU が同一のメモリを共有し、コピーを省略できる」と説明してきた。この説明は虚偽ではないが、しばしば「Apple 独自の新しいメモリアーキテクチャ」として受け取られている。実際には、UMA という手法自体は古くから存在し、Apple が実装したのは AMD が hUMA で目指した完全な共有仮想アドレス空間ではなく、その「弱い版」である。

Metal において CPU–GPU 間のコヒーレンシはコマンドバッファ境界でしか保証されず、GPU 側のデマンドページングは存在しない。共有されるのはバッファ単位の物理メモリであって、任意のポインタが両者から等価に見えるわけではない。したがって恩恵を受けられるのは、この制約に合わせて設計されたソフトウェアに限られる。

本稿は、hUMA が挫折した二つの要因(ハードウェア実装難度/プログラミングモデル変更の要求)に照らして Apple の実装を位置づけ、実際にゼロコピーを享受しているソフトウェアの実例と、享受できていない実例の双方を提示する。あわせて、共有仮想アドレス空間が提供されている他環境における利用状況を確認し、これが Apple 固有の問題ではないことを示す。

01UMA は Apple 独自の発明ではない

まず用語の整理から始める。Unified Memory Architecture という語は 1990 年代の統合グラフィックス、すなわち専用 VRAM を持たず system RAM の一部を GPU が間借りする構成に対して既に用いられていた。Intel の統合 GPU、AMD の APU、家庭用ゲーム機(PlayStation 4 / Xbox One 以降)、そして Apple 自身の iPhone / iPad 向け SoC は、いずれも物理メモリを CPU と GPU で共有する構成である。Apple Silicon の Mac は、モバイル SoC の設計をデスクトップ級に持ち込んだものであって、メモリアーキテクチャの系譜としては新規ではない。

Apple Silicon が実際に持ち込んだ差分は、アーキテクチャの新規性ではなく規模と帯域である。従来の統合 GPU における UMA は「安価だが遅い」構成の代名詞であり、ディスクリート GPU(以下 dGPU)に対する妥協として理解されていた。パッケージ内 LPDDR による広帯域と大容量の組み合わせが、この前提をコンシューマ製品で覆した点が Apple の寄与である。

A独立メモリ

CPU GPU System RAM DDR VRAM GDDR PCIe 経由のコピー

物理的に別のメモリ。転送が必ず発生する。

B従来型 UMA

CPU GPU System RAM DDR GPU 予約 帯域は CPU と共用

1990 年代の統合 GPU から存在する構成。安価だが帯域が制約。

CApple Silicon

CPU GPU NPU LPDDR 広帯域・大容量 同一パッケージ

構成は B と同型。差分は帯域と容量であって、接続の論理構造ではない。

図1 メモリ構成の三類型。B と C は論理構造として同一である。Apple Silicon が変えたのは「共有しているか否か」ではなく、共有メモリの帯域と容量が dGPU の代替たりうる水準に達したかどうかである。

しかし「物理メモリを共有している」ことと「ソフトウェアがコピーを省略できる」ことの間には、依然として大きな距離がある。この距離こそが、かつて AMD が hUMA という名で埋めようとし、埋めきれなかったものである。

02hUMA が目指したものと、挫折の構造

AMD が 2013–2014 年に hUMA(heterogeneous Uniform Memory Access)として提示し、HSA Foundation が標準化を試みた構想は、単なる物理メモリの共有ではなかった。その中核要件は次の三点である。

この三点が揃って初めて、既存の CPU 側データ構造(ポインタで連結されたツリーやリスト)を、レイアウトを変えずに GPU から辿れる。逆に言えば、hUMA の価値は「コピーが減る」ことよりも「プログラミングモデルの断絶が消える」ことにあった。

挫折の要因は二層に分かれる。ハードウェア層では、GPU のキャッシュ階層を CPU のコヒーレンシプロトコルへ参加させること、および GPU にページフォルトを扱わせることの実装コストが高い。ソフトウェア層では、たとえハードウェアが備えていても、その恩恵を受けるにはアプリケーション側が SVM を前提としたコードへ移行する必要がある。そして dGPU 環境と共存する市場では、SVM 前提のコードパスは「書いても大半のユーザーには届かない分岐」であり続けた。ソフトウェア対応は進まず、機能は使われないまま訴求点から外れていった。

補足

hUMA の系譜そのものが消滅したわけではない。AMD MI300A や NVIDIA Grace Hopper のような HPC 向け APU / スーパーチップは、CPU–GPU 間のハードウェアキャッシュコヒーレンシと GPU 側ページフォルトを実装しており、コヒーレンシの強度という一点では Apple Silicon を上回る。「一般化せずに消えた」というより、「コンシューマ市場からは退き、HPC で生き延び、Apple は別解を採った」と整理するほうが実態に近い。

03Apple が実装したもの ── 難所の解決ではなく迂回

Apple Silicon + Metal の実装を、前節の三要件に照らして検証する。

コヒーレンシの粒度

Metal の shared storage mode(MTLStorageModeShared、リソース確保時の指定としては MTLResourceStorageModeShared)は CPU と GPU の双方からアクセス可能なメモリを定義するが、Metal のヘッダに記された仕様では、コヒーレンシが保証されるのはコマンドバッファ境界においてのみである。これは CPU と GPU のキャッシュフラッシュを最小化するための設計であると明記されている。すなわち、任意のタイミングで双方が一貫した像を見る fine-grained coherency ではなく、API が定めた同期点でのみ整合する粗粒度の契約である。

fine-grained coherency ── hUMA が要件とした状態 CPU GPU 任意の時点で双方が一貫した像を見る(同期点を要求しない) Metal ── コマンドバッファ境界でのみ整合 CPU GPU commit completed 共有バッファへ書き込み 結果を読み出し GPU 実行 この区間の相互可視性は保証されない
図2 コヒーレンシが成立するタイミング。Metal の shared storage mode でも、CPU と GPU が同じ像を見ることが保証されるのはコマンドバッファ境界のみである。物理メモリは共有されているが、整合は API の契約として与えられる。

GPU 側ページフォルト

GPU がアクセスするリソースは、実行前にメモリ上へ常駐(resident)させておかなければならない。これを裏づける傍証として、Metal には MTLCommandBufferError.Code.pageFault というエラーコードが存在し、その定義は「GPU が処理できないページフォルトをコマンドバッファが発生させたことを示す」である。フォルトが解決の契機ではなく失敗として定義されている点に、設計思想が表れている。

この帰結として、Metal 4 は residency set という仕組みを導入し、ハードウェアからアクセス可能にすべきリソースをアプリケーションが明示的に指定する構造を採っている。またページ単位の柔軟な確保が必要な用途には placement sparse リソースが用意されているが、これはストレージページを持たずに確保されたリソースに対して、placement heap から アプリケーションが明示的にページを割り当てる方式であり、GPU 側のフォルト駆動ではない。要求されているのは、hUMA が自動化しようとしたメモリ管理そのものを、開発者が手作業で行うことである。

仮想アドレス空間

共有されるのはバッファ単位の物理メモリであり、CPU と GPU が同一の仮想アドレスを見ているわけではない。CPU 側ポインタは MTLBuffer.contents() を通じて得られ、GPU 側アドレスとは別に管理される。既存の malloc 済み領域を GPU に見せるには makeBuffer(bytesNoCopy:) を用いるが、この API はページアラインされたポインタを要求する。この制約自体が、行われているのがポインタの共有ではなくページテーブルのマッピングであることを示している。

hUMA ── 共有仮想アドレス空間 CPU GPU 共有仮想アドレス空間 同じポインタ値が両者でそのまま通用する Metal ── バッファ単位のマッピング CPU 仮想アドレス空間 GPU 側アドレス contents() 物理メモリ(共有 DRAM) MTLBuffer アドレス値は一致しない/マップはページ境界単位でのみ可能
図3 共有されている層の違い。Apple Silicon で共有されているのは最下層の物理メモリであり、その上の仮想アドレス空間は CPU と GPU で別に管理される。makeBuffer(bytesNoCopy:) がページアラインを要求するのは、行われているのがポインタの共有ではなくページテーブルのマッピングだからである。
表1 hUMA の要件に対する各実装の充足状況
要件hUMA / HSA 構想Apple Silicon + MetalMI300A / GH200
物理メモリの共有ありありあり
共有仮想アドレス空間フル SVMバッファ単位のマッピングSVM 対応
コヒーレンシ粒度fine-grainedコマンドバッファ境界ハードウェアコヒーレント
GPU デマンドページング必須要件非対応(エラー扱い)対応
レジデンシ管理ハードウェア/OS が解決アプリケーションが明示管理概ね自動
実際の普及限定的Apple 環境で部分的に定着HPC 領域に限定
図4 異種メモリ統合の到達度(定性評価)。0=非対応、3=完全対応として筆者が表1を数値化したものであり、測定値ではない。Apple Silicon の形状が「物理共有と実用性は高いが、抽象化の深さは浅い」ことを示す点が要点である。

整理すると、Apple は hUMA が難所とした「双方向 fine-grained コヒーレンシ」と「GPU ページフォルト」を解決したのではなく、API の契約によって回避した。同期点をコマンドバッファ境界に押し込み、レジデンシ管理を開発者に委ねることで、ハードウェア実装難度を大幅に下げている。この設計判断こそが、Apple が hUMA と同じ轍を踏まずに済んだ理由である。

04恩恵を受けているソフトウェアの実例

「メモリコピーを省略できる」という説明に対応する実装は、確かに存在する。ただしその全てが、次の三類型のいずれかに属する。

類型 A:Apple 製の専用フレームワーク(MLX)

MLX は Apple が公開した Apple Silicon 専用の配列フレームワークである。その配列は共有メモリ上に存在し、サポートされる任意のデバイス種別で、データ転送なしに演算できると明記されている。Apple 自身の説明でも、CPU–GPU 間のメモリ共有によってデータコピーが不要になり、演算は単に希望するデバイスを指定するだけであるとされる。

ここで注目すべきは API の形である。デバイスが配列の属性ではなく演算の引数になっている。これは PyTorch や NumPy の慣習からの明確な逸脱であり、共通アドレス空間を前提に書き直された専用実装であることを意味する。当然ながら Mac 専用であり、他環境でビルドすることはできない。

類型 B:クロスプラットフォーム実装内の Metal 専用分岐(llama.cpp)

llama.cpp の Metal バックエンドは、mmap された GGUF ファイルを MTLResourceStorageModeShared 経由で直接参照する。同プロジェクトの議論では、Metal 実装がメモリマップされたバッファを直接「見る」のに対し、CUDA バックエンドはデバイスバッファへコピーすると説明されている。

これが Metal 固有の特別扱いであることを示す傍証として、Intel Lunar Lake 向けの機能要望が挙げられる。そこでは、同じく UMA アーキテクチャである Apple Silicon 向けに Metal バックエンドが現在行っていることを SYCL でも概念的に模倣したい、新規デバイスバッファを確保してコピーを発行する代わりに既存のホストポインタをラップしたい、と提案されている。他バックエンドの開発者から見て「Metal だけがやっている特殊なこと」と認識されているわけである。

同時に、この方式の限界も実装上の不具合として露出している。Metal での部分オフロードは、そのバックエンドに割り当てられた全テンソルにまたがる単一の mmap 領域としてウェイトをマップするため、テンソルがファイル内で連続配置されていない GGUF では過剰なマッピングが発生する。報告例では、49 層中 2 層(約 947 MiB)しかオフロードしない設定で、モデルファイル全体の約 30 GB がマップされ、GPU のワーキングセット上限を超過している。バッファ単位・事前レジデンシという制約が、そのまま実用上の副作用として現れた事例である。

類型 C:元から Apple 専用の API(IOSurface 系)

IOSurface に裏打ちされた CVPixelBuffer を CVMetalTextureCache 経由で Metal テクスチャとして参照し、AVFoundation / VideoToolbox / Core Image の間をコピーなしで繋ぐ映像パイプラインは、macOS / iOS で最も広く使われている共有メモリの実例である。ただしこれは Apple Silicon 以前から存在する抽象化であり、Intel Mac + dGPU の時代には内部でコピーが発生していた。UMA が生んだ恩恵というより、UMA によって初めて宣伝どおりに機能するようになった既存の抽象化と位置づけるべきである。

05恩恵を受けていない実例 ── PyTorch MPS

前節の対照として、最も示唆的なのが PyTorch の MPS バックエンドである。2026 年 1 月に提出された機能提案によれば、現在の MPS バックエンドではテンソルのバッキングバッファは所有デバイスのプライベートメモリに確保され、CPU と MPS の間でテンソルを移動すると新しいバッキングバッファが確保される。ユニファイドメモリを活用して読み取り専用テンソルの memcpy を回避し、所有権の変更のみで済ませることでアロケーションを半減させる案は、この時点でまだ提案段階である。

さらに重要なのは、その提案自体が可変テンソルについては断念している点である。一方への書き込みが他方を変えてはならないという PyTorch のテンソル所有権セマンティクスを維持する以上、ソースとデスティネーションで二つのコピーを保持せざるを得ないと結論づけられている。

すなわち、ハードウェアが物理コピーを不要にしても、フレームワークの意味論が論理的にコピーを要求する。M1 の登場から五年以上が経過し、Mac が LLM のローカル実行環境として一定の地位を占めた後でさえ、支配的なクロスプラットフォームフレームワークではゼロコピー化が未完了である。hUMA においてソフトウェア対応が進まなかった構造が、そのまま再現されている。

A CUDA + dGPU GGUF ホストバッファ デバイスバッファ GPU mmap cudaMemcpy(PCIe) B llama.cpp / Metal GGUF mmap したページ=MTLBuffer GPU mmap 直接参照(storageModeShared) C PyTorch / MPS CPU テンソル MPS テンソル GPU memcpy 同一の物理メモリ内での複製
図5 ウェイト/テンソルがたどる経路の比較。B ではファイルのページがそのまま GPU から見えるのに対し、C では物理的な転送が不要であるにもかかわらず、フレームワークの所有権セマンティクスによってコピーが発生する。ハードウェアが可能にしたことと、ソフトウェアが実際に行っていることの差はここに現れる。
表2 主要ソフトウェアにおけるゼロコピー活用状況
ソフトウェア類型ゼロコピー他環境への移植性
MLXApple 製専用全面的なし(Mac 専用)
llama.cpp (Metal)専用分岐ウェイトのみ本体は移植可、当該パスは Metal 限定
IOSurface 系パイプラインApple 専用 API全面的なし
PyTorch (MPS)移植未対応(提案段階)高い

06SVM が実在する環境での利用状況

ここまでは Apple 環境の内部で議論を閉じてきたが、それでは「Apple が特に劣っている」という誤読を招く。共有仮想アドレス空間(SVM)が実際に提供されている環境において、それがどの程度使われてきたかを確認しておく必要がある。

OpenCL 2.0 SVM の顛末

OpenCL 2.0 は 2013 年に SVM を導入した。ポインタで連結されたデータ構造をそのままカーネル引数として渡せること ── 第 2 節で見た hUMA の中核要件と同じもの ── が、その最大の利点として掲げられた。

しかし仕様の構造そのものに問題があった。全 OpenCL 2.0 プラットフォームに実装が必須とされたのは coarse-grained buffer SVM のみで、fine-grained buffer SVM と fine-grained system SVM は任意扱いである。すなわち、hUMA が本来目指した粒度は仕様上の必須要件から外された。

実装側の反応はさらに厳しかった。分子動力学の OSS である OpenMM が fine-grained buffer SVM を実装した際の記録によれば、Intel と AMD は 2014 年に OpenCL 2.0 対応ドライバを速やかに提供したが、NVIDIA は評価段階を超えず、Apple は OpenCL 1.2 以降のサポートを停止した。結果としてこの最適化は Linux と Windows 上の AMD / Intel GPU に対してのみ意味を持ち、Mac で OpenCL 2.0 を回避できるよう USE_OPENCL_120 の define によって切り分ける構成が採られた。

性能は良好であったと報告されている。それでもなお、単一のコードベースを二系統に分岐させ、その両方を保守する形でしか出荷できなかった。クロスプラットフォームな OSS における SVM の帰結が、この一例に集約されている。

単一の コードベース #ifdef 従来経路(明示コピー) 全プラットフォーム NVIDIA・Apple を含む SVM 経路 AMD・Intel Linux/Windows 対象外 同一機能について二系統を保守することになる
図6 クロスプラットフォームな OSS が SVM を採用したときに生じる構造。性能上の利益があっても、対象範囲が狭い経路のために保守対象が二重化する。合理的な判断として、多くのプロジェクトはこの分岐を書かない。

同型の痕跡は llama.cpp にも残っている。CUDA バックエンドには GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 という環境変数が用意され、Linux 上でユニファイドメモリを有効化できるとされている。既定では無効であり、対象 OS も限定されている。中核の実行パスには置かれていない。

例外としての HPC

明確な例外が HPC 領域に存在する。AMD は MI300A について、単一メモリ空間がデータ複製と転送を不要にし、アプリケーションの段階的な高速化を可能にするとした上で、その潜在能力は OpenMP 5.2 の抽象化 ── ホストとデバイスのデータ環境を統合できる ── を通じて実現できると整理している。事例として提示されたのが OpenFOAM である。移植性の高いオープンソースの C++ CFD ライブラリを、本番運用規模のまま、ディレクティブベースの OpenMP でオフロードしている。

ここでは SVM が、移植性を損なう分岐としてではなく、移植コストを下げる手段として機能している。CPU と GPU が同一の物理ページに対してロード/ストアを発行するため、アプリケーションチームはユニファイドメモリのプラグマを挿入する漸進的な移植戦略を採ることができる。既存の巨大な CPU コードベースを、書き換えずに段階的に GPU 化したいという要求に対して、SVM は正面から応えている。

表3 共有アドレス空間の提供状況と実利用
環境提供されるもの汎用 OSS での実利用
OpenCL 2.0 SVMcoarse-grained は必須、fine-grained は任意分岐前提。NVIDIA・Apple が対象外
CUDA(llama.cpp の例)ユニファイドメモリのオプトイン既定無効・Linux 限定
OpenMP 5.2 + MI300Aプラグマによる統合データ環境OpenFOAM 等で実利用(HPC 限定)
Metalバッファ単位の物理メモリ共有のみllama.cpp 等で実利用(SVM ではない)

成立条件を個々のソフトウェアに適用する

この例外を成立させている条件は、次の三点に分解できる。

これを「HPC 対 コンシューマ」という軸で適用するのは誤りである。Apple のプラットフォームは、①と②をむしろ HPC より強く満たしているからである。

①について、HPC では調達単位でハードウェアが確定するのに対し、Apple ではプラットフォームの定義として確定している。Apple Silicon 上で GPU を用いる限り、UMA でない構成は存在しない。②についても、Clang / LLVM への長期の投資、Xcode、Metal のシェーダコンパイラ、Accelerate、Core ML、そして MLX に至るまで、ツールチェーンとライブラリの全層を単一の主体が供給している。さらに Apple は、PyTorch の MPS バックエンドや TensorFlow の Metal プラグインのように、他者のフレームワークへの実装までを自ら供給している。ベンダとラボが一体で移植を担うという HPC の構図と、ほぼ同型である。

分岐するのは③のみである。そしてこれはプラットフォームの属性ではなく、個々のソフトウェアの属性である。同じ Apple Silicon 上で、同じベンダのエコシステムを用いていても、③を満たすかどうかはソフトウェアごとに異なる。

表4 三条件の充足と、第 4 節で見たゼロコピー達成度の対応
ソフトウェア①②③ゼロコピー
MLX○○○全面的
IOSurface 系パイプライン○○○全面的
llama.cpp (Metal)○○△ウェイトのみ
PyTorch (MPS)○○×未対応

①②はプラットフォームの属性であり、Apple Silicon 上のソフトウェアには一律に与えられる。③のみがソフトウェアごとに異なる。llama.cpp の △ は、本体の移植性を保ったまま Metal 専用の分岐を持つ範囲でのみ満たすことを示す。

この枠組みをソフトウェアの類型に一般化すると、Apple 専用として書かれたサードパーティ製アプリケーションは③を満たす側に、複数 OS を対象とする汎用 OSS は満たさない側に置かれる。ただし後者の内部には幅がある。llama.cpp と PyTorch はいずれも汎用 OSS でありながら結果が分かれており、分かれ目は「他環境を対象とするか否か」ではなく、環境固有の分岐をどの層まで許容する設計かにある。

なお、Apple 専用の商用アプリケーションについては、外部から実装を確認する手段がないため、本稿では類型として述べるに留める。表 4 に挙げた実装は、いずれも公開された一次資料によって確認できるものに限っている。

①と②が共通である以上、結果を分けているのは③だけである。第 4 節で三類型として分類したものは、この条件から説明できる。すなわち三条件は事後的な分類ではなく、どのソフトウェアがゼロコピーを回収しうるかを事前に判定する枠組みとして機能する。

Windows PC と Android については、事情が異なる。dGPU と iGPU が混在し、SoC のベンダも複数存在するため①を満たさず、エコシステムを構築する単一の主体も存在しないため②も満たさない。三条件のうち脱落する位置が、Apple とは根本的に異なる。

もう一つの障壁 ── 抽象化レイヤとの衝突

ハードウェアの可用性とは別に、フレームワーク側にも構造的な障壁がある。llama.cpp の ggml_backend も PyTorch の device も、「バックエンドはメモリを所有し、境界で転送を行う主体である」というインターフェースとして設計が固まっている。SVM はこの層そのものを不要にするため、恩恵を受けるにはバックエンドを一つ追加するのでは足りず、フレームワークの中核的な意味論を変更する必要がある。

第 5 節で見た PyTorch MPS の限界 ── 可変テンソルのゼロコピー化がテンソル所有権セマンティクスによって阻まれる ── は、まさにこの層の問題である。Metal のバッファ単位共有が「バックエンドを一つ書く」で到達できるのに対し、SVM は「フレームワークの意味論を変える」まで到達しなければ利用できない。難度には桁の差がある。OpenFOAM が越えられたのは、OpenMP というディレクティブの抽象化層に unified memory の概念が標準として後から追加されたためであり、フレームワーク側が対応した稀な例と見るべきである。

なお、SVM が利用可能であることと高速であることは別である。ページフォルトとページマイグレーションのコストがある以上、最適化を突き詰めた実装ほど明示的なコピーへ回帰する傾向がある。

なぜ Apple は SVM を実装しなかったのか

本項は筆者の解釈である

以下は前項までの事実から導いた推論であり、Apple の設計判断に関する一次資料に基づくものではない。

三条件のうち①と②を HPC 以上に満たしているならば、Apple には hUMA 相当を実装する条件が揃っていたことになる。それをしなかったことには説明が要る。

Apple は HPC と同じ条件を備えながら、逆の設計判断を採った。HPC は強い共有機構をハードウェアに実装し、少数のアプリケーションを個別に手作業で移植する。Apple は弱い共有機構に留め、多数のアプリケーションが特別な努力なしに恩恵を受ける構造を選んだ。

反転の理由は、対象となるアプリケーション数の桁にあると考える。El Capitan で移植対象となるコードは数十本の規模であり、一本あたりに人年単位の費用を投じても引き合う。Apple が相手にするサードパーティは数百万本であり、一本あたりの追加コストがゼロに近くなければ普及しない。同じ垂直統合であっても、対象数の桁が四つも五つも異なれば最適解は反転する。

もう一点、垂直統合が届く範囲の差がある。HPC では書き換える対象がラボ自身のコードであり、抽象化層である OpenMP も標準化の過程を通じてベンダの影響下にある。両端に手が届く。Apple はエコシステムの全層を供給しながら、前項で見たとおり PyTorch の中核的な意味論には手を出せなかった。自らが所有していない抽象化は、書き換えられない。第 5 節で観察した限界は、Apple の技術力ではなく、垂直統合の及ぶ範囲の境界に由来している。

整理すると、共有仮想アドレス空間の恩恵は、三条件を満たす個々のソフトウェアにおいてのみ回収されている。ハードウェアが確定し、単一の主体がエコシステムを供給していても、他環境での動作を要件に含むソフトウェアは、抽象化レイヤの壁によってそれを回収できない。これは Apple 固有の問題ではなく、異種メモリ統合という課題に共通する構造である。

07なぜ Apple では部分的に定着したのか

本節は筆者の解釈である

以下は前節までの事実から導いた推論であり、Apple や AMD の意思決定に関する一次資料に基づくものではない。

前節では、Apple が SVM を実装しなかったことの合理性を論じた。本節で扱うのは逆の問いである ── より弱い共有機構が、なぜ実際に使われたのか。技術的には、Apple の実装は hUMA より弱い。にもかかわらず Apple 環境でのみ一定の定着を見た理由は、思想の優劣ではなく垂直統合が生んだ非対称性にあると考える。

AMD にとって、hUMA 対応コードパスは常に「オプションの最適化」であった。同一のバイナリが dGPU 環境でも動作しなければならず、SVM を前提とした実装は分岐として追加され、テストされ、保守される必要があった。投じたコストに対して届くユーザーは APU 利用者に限られる。合理的な開発者は、この分岐を書かない。

Apple Silicon では、Metal バックエンドを書くことが「唯一のパス」である。Mac 上で GPU を使う以上、実質的に Metal しか選択肢がなく、そしてその Metal が動く全てのハードウェアが UMA である。分岐が生まれるのは Metal 対 CUDA という境界であって、UMA 対 dGPU という境界ではない。結果として、Metal バックエンドを書いた時点で自動的に UMA 前提のコードになる。llama.cpp の Metal バックエンドがゼロコピーであるのは、開発者が UMA のために特別な努力をしたからというより、Metal で書けばそれが最も自然な実装だったからだと見るのが妥当である。

もう一点、Apple の設計が「弱い版」であったこと自体が普及に寄与している。fine-grained SVM への対応は既存コードの深い書き換えを要求するが、「バッファを shared storage mode で確保し、コマンドバッファ境界で同期する」という契約は、既存のグラフィックス API の作法とほぼ同じである。移行コストが低いところで妥協した設計が、結果として採用された。

08結論

hUMA が答えようとした問い ── 異種プロセッサ間でプログラミングモデルの断絶をどう消すか ── に対して、Apple は答えを出していない。断絶を消す代わりに、断絶の位置を動かしにくい場所(コマンドバッファ境界)に固定し、そこを跨ぐコストをゼロに近づけた。これは工学的には賢明な妥協であり、マーケティング上は「統合」と呼びうるものである。だが hUMA が目指した地点とは異なる場所にあることは、記録しておく価値がある。

09補講 MacBook はなぜ 8GB で動くのか

本章は本論から外れる。ユニファイドメモリは、しばしば「メモリが少なくても済む」ことの理由として説明される。この説明の当否は本稿の主題ではないが、UMA に対する誤解として最も広く流通しているものであり、これまでの整理を適用すれば判定できる。

現行ラインナップにおける 8GB の位置

まず前提を確認する。Apple は 2024 年 10 月 30 日に MacBook Air の基本構成を 8GB から 16GB へ変更した。米国での開始価格は据え置かれている。M5 MacBook Air も全モデルが 16GB 標準であり、2026 年 3 月のリフレッシュ以降、MacBook の全構成が 16GB 始まりとなった。

唯一の例外が MacBook Neo である。2026 年 3 月発売の最廉価機で、8GB の LPDDR5X を搭載し、増設オプションを持たない。ただしこの 8GB は「UMA だから足りる」という設計判断の結果ではない。Neo は iPhone 16 Pro の A18 Pro をそのまま流用しており、この SoC は TSMC の InFO-PoP パッケージによって 8GB の LPDDR5X をダイ直上に積んだ状態で出荷される。100 mm² を超える基板面積を節約する代わりに、構成は SoC が元々パッケージされている容量に縛られる。物理的な上限であって、性能上の見積もりではない。 後継機では A19 Pro とともに 12GB へ増えると報じられている。

UMA が実際にバイト数を減らす経路

効果がないわけではない。UMA が実メモリ消費を減らす経路は三つある。

いずれも実在する効果である。しかし節約されるのはGPU やメディアエンジンに隣接したデータに限られる。8GB のマシンを逼迫させるのは、まずそれではない。ブラウザのタブ、Electron ベースのアプリケーション、統合開発環境、写真ライブラリのインデックス ── いずれも CPU 側のメモリ圧である。恩恵のカテゴリと逼迫のカテゴリが一致していない。

ただし一つだけ、宣伝されないまま常時働いている経路がある。ウィンドウのバッキングストアは IOSurface で裏打ちされており、それを合成する WindowServer は Apple 自身が実装している。Apple Silicon の UMA には宣伝されない実利がある ── GPU 向けの領域を起動時に切り出さず、必要になった時点で確保し不要になれば解放するため、確保量が実需要に追従する。GPU 側の需要が小さい間はその分がそのまま CPU 側で使え、メモリの利用効率を高めやすい。加えてコンポジタは両端を Apple が握っているため、サードパーティの対応を一切要さずに CPU / GPU 間のコピーを回避している。第 4 節で見た類型のうち、ゼロコピーが全ユーザーに対して普遍的に成立している数少ない領域である。この効率が、MacBook が 8GB 構成でも軽作業であれば快適に動作する一因である可能性は高い。

切り出しの有無 ── Windows との比較

Windows 環境では、GPU 向けの切り出し量を変更するのに再起動を要する。AMD の Variable Graphics Memory も Intel の Shared GPU Memory Override も、ドライバの GUI から設定できるようになっただけで、反映は再起動後である。Windows にも動的に割り当てられる共有 GPU メモリ層は存在するが、多くのアプリケーションが dedicated video memory の報告値で挙動を決めるため、大容量を要する用途では固定的な切り出しを設定せざるを得ない。その設定はピーク需要に合わせることになり、差分は使われないまま OS から見えなくなる。

macOS には切り出しそのものが存在しない。recommendedMaxWorkingSetSize という上限値があるだけで、上限は消費を伴わない。実際に確保されるのはバッファ単位の要求があった時点であり、解放されればそのまま CPU 側に戻る。動かしているのは境界ではなく閾値である。

Windows(切り出しあり) 切り出し(ブート時に固定) 時間 → GPU 側遊休は CPU 側からは使えない macOS(切り出しなし) GPU 確保量の上限値(消費を伴わない) 時間 → 空きは CPU / GPU いずれも使える GPU 実使用 CPU 実使用 GPU 遊休(Windows のみ) 空き
図7 同じ需要の推移に対する挙動の違い。両パネルの GPU 需要と CPU 需要は同一である。Windows では切り出しをピーク需要に合わせるため、需要が小さい間の差分が遊休となり、その分は CPU 側からは使えないまま OS の管理外に置かれる。macOS は切り出しを持たず、GPU も CPU も実需要に追従して確保・解放するため、余りが分断されず一つの空き領域になる。上限値は GPU の確保量だけに掛かるもので、CPU 側の使用量とは無関係であり、それ自体は物理メモリを消費しない。

統合プールは小容量機では不利に働く

さらに、統合プールであること自体が小容量構成では逆向きに作用する。

dGPU 搭載 PC System RAM 8GB CPU が専有 VRAM 8GB GPU が専有 合計 16GB Apple Silicon(UMA) ユニファイドメモリ 8GB CPU と GPU が奪い合う GPU 作業セット上限(搭載量の一部) 合計 8GB 重複排除による節約は、まずこの差を埋めるために費やされる
図8 同じ「8GB」の意味の違い。dGPU 機の 8GB は CPU が専有できる量であるのに対し、UMA の 8GB は CPU と GPU の作業セットが共有する総量である。加えて Metal が GPU に許す作業セットは搭載量の一部に制限されるため、小容量構成では GPU 側の上限自体も低くなる。

すなわち、8GB の UMA と「8GB の RAM に加えて専用 VRAM を持つ構成」は等価ではない。重複排除による節約は、この構造的な差を埋め合わせる方向にまず消費される。差し引きで有利になるかどうかは GPU 側の作業量に依存し、GPU をほとんど使わない用途では節約分もゼロである。

8GB を成立させている実際の要因

本項は筆者の解釈である

以下の要因の寄与度について、Apple による公式の内訳は示されていない。

8GB の Mac を実用範囲に収めてきたのは、前項で述べたコンポジタ経路を別にすれば、UMA ではなく次の三つであると考える。メモリ圧縮 ── macOS は非活性ページを積極的に圧縮する。高速な内蔵 SSD へのスワップ ── 実効的なメモリ階層の延長として働く。アプリケーション側の実装 ── ネイティブアプリの多くが比較的軽量に書かれている。

いずれも UMA とは独立した機構である。垂直統合の成果ではあるが、CPU と GPU がメモリを共有していることとは関係がない。にもかかわらず「8GB で足りる」の説明に UMA が持ち出されるのは、共有という事実が効率的に聞こえるためであろう。

用途別に見た実効性

第 4 節までの整理を用途に適用すると、UMA の恩恵は三つの軸に分解できる。①容量プーリング(GPU が dGPU の VRAM 容量を超える対象を扱える)、②転送量(CPU と GPU の間を往復する総バイト数)、③往復頻度(一回は小さくても交互実行が高頻度)である。恩恵の源泉が異なるため、8GB 構成での実効性も異なる。

表5 用途別に見た UMA の恩恵と、8GB 構成での実効性
用途主たる軸8GB 構成での実効性
オフラインレンダリング①反転する。容量が前提の恩恵であり、小容量では成立しない
映像編集・エンコード③有効。メディアエンジンと GPU と CPU が同一メモリを見る
ローカル LLM 推論①最も不利。dGPU 機に対する優位が消える
リアルタイム 3D③有効だが、作業セットが収まる範囲に限られる
一般的な事務・ブラウジング③アプリケーション側には恩恵が生じない。OS のコンポジタ経路のみが常時働くが、支配的なのは CPU 側のメモリ圧である

8GB の Neo が実際に得ている恩恵は、このうち③に集中する。Neo のメディアエンジンは H.264 / HEVC / ProRes / ProRes RAW のハードウェアアクセラレーションと AV1 デコードを備え、メモリ帯域は 60GB/s である。軽い映像処理においてコピーが消えることは事実である。一方で、より重いクリエイティブ用途では上位機に及ばないと評価されており、Apple 自身も動画編集やデータ分析を Neo の想定用途から外している。

小結

以上から、次のように整理できる。

UMA が正当に主張できるのは「同じ作業を dGPU 機と比べたとき、DRAM と VRAM の合計は少なくて済む」ことであって、「単一プールの容量が小さくて済む」ことではない。「メモリが少なくても済む」という表現は通常後者と解釈されるため、説明として適切ではない。

この判定は、本論の結論と同じ構造をしている。UMA は物理メモリの共有を実現したが、それによって何が得られるかは用途とソフトウェアの設計に依存する。共有という事実から、容量の節約も、コピーの省略も、自動的には導かれない。

10検証の手立て

本稿の主張は、以下の実測によって手元で確認できる。

11出典

  1. Metal.framework MTLResource.h ── storage mode ごとのコヒーレンシ保証範囲の記述。 https://github.com/xybp888/iOS-SDKs/blob/master/iPhoneOS13.0.sdk/System/Library/Frameworks/Metal.framework/Headers/MTLResource.h
  2. Apple Developer Documentation ── MTLCommandBufferError.Code.pageFault。 https://developer.apple.com/documentation/metal/mtlcommandbuffererror/code/pagefault
  3. WWDC25 "Discover Metal 4" ── residency set および placement sparse リソース。 https://developer.apple.com/videos/play/wwdc2025/205/
  4. Metal Feature Set Tables ── supportsPlacementSparse の対応世代。 https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf
  5. ml-explore/mlx ── 共有メモリ上の配列と転送なしのデバイス間演算。 https://github.com/ml-explore/mlx
  6. WWDC25 "Get started with MLX for Apple silicon"。 https://developer.apple.com/videos/play/wwdc2025/315/
  7. ggml-org/llama.cpp Discussion #21223 ── Metal は mmap バッファを直接参照、CUDA はコピー。 https://github.com/ggml-org/llama.cpp/discussions/21223
  8. ggml-org/llama.cpp Issue #21827 ── SYCL バックエンドで Metal 同様のゼロコピーを求める要望。 https://github.com/ggml-org/llama.cpp/issues/21827
  9. ggml-org/llama.cpp Issue #24510 ── Metal での部分オフロード時の過剰マッピング。 https://github.com/ggml-org/llama.cpp/issues/24510
  10. pytorch/pytorch Issue #172987 ── MPS バックエンドにおけるユニファイドメモリ活用の提案(2026-01)。 https://github.com/pytorch/pytorch/issues/172987
  11. openmm/openmm Issue #3947 ── OpenCL 2.0 fine-grained SVM の実装と、ベンダ対応状況による分岐。 https://github.com/openmm/openmm/issues/3947
  12. Intel "OpenCL 2.0 Shared Virtual Memory Overview" ── SVM の三階層と必須/任意の区分。 https://www.intel.com/content/www/us/en/developer/articles/technical/opencl-20-shared-virtual-memory-overview.html
  13. ggml-org/llama.cpp docs/build.md ── GGML_CUDA_ENABLE_UNIFIED_MEMORY の位置づけ。 https://github.com/ggml-org/llama.cpp/blob/master/docs/build.md
  14. Tandon et al., "Porting HPC Applications to AMD Instinct MI300A Using Unified Memory and OpenMP" (arXiv:2405.00436) ── OpenFOAM の事例。 https://arxiv.org/abs/2405.00436
  15. Apple Support ── MacBook Air (13-inch, M5) Tech Specs(16GB 標準構成)。 https://support.apple.com/en-us/126320
  16. Apple Support ── MacBook Neo (13-inch, A18 Pro) Tech Specs(8GB、メディアエンジン、60GB/s)。 https://support.apple.com/en-us/126322
  17. MacRumors ── MacBook Air 基本構成の 8GB から 16GB への変更(2024-10-30)、および MacBook Neo の 8GB 固定。 https://www.macrumors.com/2026/03/04/macbook-neo-8gb-ram/
  18. TechPowerUp ── A18 Pro の InFO-PoP パッケージによる 8GB の物理的制約。 https://www.techpowerup.com/347063/apple-macbook-neo-capped-at-8-gb-ram-by-a18-pro-info-pop-packaging
  19. Intel Support ── Shared GPU Memory Override の要件。既定値 57%、反映には再起動を要する。 https://www.intel.com/content/www/us/en/support/articles/000101789/graphics.html
  20. XMG Help Center ── AMD Adrenalin の Variable Graphics Memory による iGPU 向けメモリ割り当て。BIOS 設定は Auto のまま、ドライバ側で指定する。 https://help.xmg.gg/hc/en-gb/articles/30052785697437-Allocation-of-video-memory-in-systems-with-integrated-graphics-units-iGPU