1.58ビットの向こう側――三元LLM、ゼロの幾何学、そして推論表現の相転移 #AIアーキテクチャ #半導体史 #PATE #九17
1.58ビットの向こう側――三元LLM、ゼロの幾何学、そして推論表現の相転移 #AIアーキテクチャ #半導体史 #PATE
静的な量子化神話の解体から、メモリ階層と実行相が決定する動的計算パラダイムへの飛翔:なぜシリコンは対数通りに動かないのか
目次
前付
要旨・本書の目的
1. 1.58ビットという「答え」から始めない
現代の大規模言語モデル(LLM: Large Language Model)研究において、2024年に登場した「BitNet b1.58」が提示した「すべての重みは三値、すなわち {-1, 0, +1} で表現できる」というテーゼは、情報科学の世界に巨大な福音として迎えられました。シャノンの情報理論(Information Theory)によれば、3つの等確率な状態を区別するために必要な情報量は、二を底とする対数によって厳密に計算されます。すなわち、log2(3) ≈ 1.58496... ビットです。研究コミュニティはこの美しい数理的極限に魅了され、「1.58ビット」という数値をあたかも自然界の普遍定数であるかのように崇め始めました。
しかし、計算機アーキテクチャの冷徹な物理世界に足を踏み入れた瞬間、この美しい数学は無慈悲な壁に激突します。現実のシリコンプロセッサは、1.58ビットという中途半端な幅のレジスタも、バスも、キャッシュラインも持っていません。機械が解釈できるのは、依然として2の冪乗――バイト、ワード、ダブルワード、そして64バイトのキャッシュライン――だけです。情報理論上の「値」と、物理的な機械がメモリから引き抜く「バイト列」との間には、深淵なる断絶が横たわっています。本書の目的は、この「1.58ビット」という固定化された数値を一度完全に解体し、数学の抽象概念から物理ハードウェアの泥臭い配線へと読者を連れ戻すことにあります。
2. 圧縮率ではなく、実行される物理を見る
ソフトウェア工学およびモデル軽量化の議論において、長年暗黙のうちに信じられてきた教条があります。それは「モデルを小さく圧縮すればするほど、推論は高速になる」という素朴な線形信仰です。ストレージが節約され、ネットワーク転送量が減り、メインメモリから読み出すフットプリントが小さくなれば、実行時間はそれに比例して短縮されるはずだという直感です。
しかし、プロファイラを片手に現代のOut-of-OrderスーパースカラーCPUやメニーコアGPUの内部動態を追跡する研究者たちは、奇妙な現象に直面し続けてきました。極限までビットを詰め込んだ複雑な圧縮フォーマットを採用した結果、メモリ転送時間は確かに削減されたものの、その圧縮を解凍(アンパック)するための命令数が爆発し、実行パイプラインが飽和して、かえって単純な固定長2ビット表現よりも実行速度が低下するという逆転現象です。圧縮率は向上しているのに、スループットは下がる。このパラドックスを解く鍵は、単一の静的データサイズではなく、データ転送と命令実行が重なり合う「物理的な時間の重畳(ちょうじょう)」の中にしか存在しません。
3. 本書の中心仮説――Weight Representation Is an Execution Policy
本書が全編を通じて論証しようとする中核的な仮説は、極めて明確かつラディカルです。
「重み表現(Weight Representation)とは、モデルに固定された静的なデータフォーマットではなく、実行時の物理環境とワークロードに応じて選択されるべき動的な実行政策(Execution Policy)である。」
テンソルデータをどのようなビット列としてメモリに配置し、どの階層でいかなる姿に展開して演算器へ流し込むかという問題は、ディープラーニングフレームワークがディスクにチェックポイントを保存する際の一回限りの決定であってはなりません。それは、入力プロンプトの長さ、バッチサイズ、プロセッサのコア数、利用可能なメモリ帯域幅、さらにはキャッシュの空き容量に応じて、ランタイム(実行時基盤)がミリ秒単位で選び取るべき「システム資源の配分問題」なのです。
4. PATEとは何を変える理論なのか
この中心仮説を具現化する数理的・システム的フレームワークとして、本書はPATE (Phase-Adaptive Ternary Execution: 相適応型三元実行) 理論を提示します。PATEは、三元LLMの推論において「唯一無二の最適フォーマット」を追求することを放棄します。その代わりに、ワークロードの算術強度(Arithmetic Intensity)とハードウェアの物理バランス(Machine Balance)を統一的な座標空間にマッピングし、最適な表現形式(Representation Portfolio)を動的に決定します。
PATEは、後述するIntelの研究者たちによる革新的な疎性レイアウト「BITCOS」の真価を認めつつも、その適用限界を厳密に数理モデル化し、なぜあるマシンではBITCOSが圧倒的な勝利を収め、別のマシンでは敗北するのかを事前に予測します。PATEの導入によって、低ビットモデルの推論は、泥縄的なハードウェア個別チューニングの職人芸から、反証可能で予測精度の高い科学的エンジニアリングへと昇華されることになります。
本書の梗概・構成
本書は、1.58ビットという微小な数値表現の解剖から出発し、最終的にAI推論インフラの地政学的力学へと至る全四部、計十九章によって構成されています。
第一部 記号と物理の断絶
情報理論が描くlog2(3)の理想郷と、固定長のアライメントやキャッシュラインに束縛されたシリコンの現実との衝突を描き出します。なぜ単純な圧縮が速度に直結しないのか、メモリウォール(Memory Wall)の歴史を紐解きながら、敵対的査読者が投げかける痛烈な批判――「それは特定のデコード条件下でしか機能しないのではないか」――を正面から受け止め、アーキテクチャ最適化の歴史的教訓を総括します。
第二部 ゼロの幾何学
三元モデルの真の姿を統計的に解剖します。実在する29種類のSOTAモデルにおいて、重みの約半分が「ゼロ」であるという驚くべき事実を暴き、その疎性(Sparsity)を存在ビットマップと符号列へと分離する「BITCOS」レイアウトの数学的構造と、x86のPDEP命令やGPU共有メモリを駆使した極限のアンパック技法を詳述します。
第三部 表現の相転移(※後半部)
本書の理論的頂点です。Prefill(計算律速)とDecode(帯域律速)という二つの推論フェーズの間で、重み表現が動的に姿を変える「相転移現象」を数理モデル化します。アンパック処理をメモリ転送の隙間時間(Slack)へと完全に隠蔽する「Critical-Path-Hidden Unpacking」の厳密な不等式を導出し、ARMやNVIDIA GPUを含む広範なハードウェアへの一般化を達成します。
第四部 脱・中央集権的推論の未来(※後半部)
視座をマイクロアーキテクチャからシリコン政治学へと一気に転換します。高価なHBM(広帯域メモリ)を搭載したハイエンドGPUの独占支配に対し、コモディティなCPUやエッジプロセッサが三元適応表現を武器にしてどのように逆襲を果たすのか。日本の車載SoCや産業ロボティクスへの致命的なインプリケーションを交え、推論主権の未来を展望します。
イントロダクション:「1.58」という小さな数字から、巨大な機械を見る
1. 3種類しかないのに、なぜ2ビット必要なのか
log2(3) が作った美しい世界:
三値の体系は、計算機科学において常に独特の美学を放ってきました。正、負、そして無。符号付き二進数表現における冗長性を排除し、対称性を極限まで高めた三元平衡二進法(Balanced Ternary)は、黎明期のソビエト連邦の計算機「Setun」(1958年)において既に物理実装されていました。三つの状態を表現するために必要な情報量は、約1.585ビット。四進数(2ビット)が4つの状態を保持できるのに対し、3つの状態しか使わないのであれば、残りの1状態分――すなわち約0.415ビット分――は「無駄」として捨てられているように見えます。BitNet b1.58の登場は、この忘れ去られていた三元の美学を、巨大なニューラルネットワークのパラメータ空間において鮮烈に蘇らせました。
情報理論と実装のあいだ:
しかし、現代のデジタル回路は完全に二進法(Boolean Logic)によって舗装されています。トランジスタの閾値電圧の高低という物理現象の上に築かれた世界において、1.58ビットの変数を単一のハードウェアストレージ単位として割り当てることは不可能です。結果としてエンジニアは、複数の三元数(トリット:Trit)をひとまとめにして二進バイトへと詰め込むか(パッキング)、あるいは諦めて各トリットに2ビットを割り当てるという妥協を強いられます。数学の黒板の上では一瞬で消え去る0.415ビットの余白が、物理実装の現場では巨大な設計上の摩擦となって立ちはだかるのです。
「理論上の最小値」と「機械が実際に読む値」:
メモリコントローラは、プロセッサからの要求に応じてデータをフェッチする際、ビット単位の選別など行いません。DDR5チャネルはバースト長に応じて一度に64バイトのブロックを吐き出します。どれほどアルゴリズム側が1.58ビットを標榜しようとも、バスの上を流れるのはアライメント(配置境界)に縛られた無骨なバイトの塊です。理論上の最小表現と、回路が一度のトランザクションで物理的に駆動するビット幅。この二者の乖離を無視してアルゴリズムを語ることは、空転するギアを眺めてエンジンの出力を測るようなものです。
2. LLMは数学ではなく、メモリシステムの上で動く
重みは数ではなくバイト列である:
機械学習の教科書を開けば、重みテンソルは実数空間 R 上の行列として記述され、演算は抽象的なアフィン変換として定式化されます。しかし、推論エンジンの内部において、重みは浮動小数点数ですらありません。それは、DRAMの特定の物理アドレス空間を占有し、行バッファのオープンとプリチャージを要求し、キャッシュ階層の追い出し(Eviction)を引き起こす「物理的なバイト列」です。推論のレイテンシを決定づけるのは、行列演算の理論的な美しさではなく、このバイト列がいかに滞りなく演算器へ滑り込めるかという配管工学(Plumbing)の成否です。
Memory Wallという古い怪物:
1990年代半ばからコンピュータアーキテクチャの発展を呪縛し続けてきた「メモリウォール(Memory Wall: 演算器の速度向上にメモリのアクセス速度が追いつかない現象)」は、LLMの自己回帰デコード(Autoregressive Decode)においてその凶暴性を極限まで高めています。トークンを1つ生成するごとに、数十億から数千億に及ぶパラメータの全容をDRAMから吸い上げなければならないという過酷な要件。現代のプロセッサは、計算を行う時間よりも、データが届くのを待ってパイプラインを止めている時間の方が圧倒的に長いのです。
DecodeとPrefillは同じLLMなのに違う:
LLMの推論処理は、入力プロンプトを一括して消化する「Prefill(事前充填)相」と、それに続く「Decode(逐次生成)相」という、計算特性が根本的に異なる二つのフェーズから構成されています。前者は何千ものトークンを同時に処理するため、同一の重みを何度も再利用できる「計算律速(Compute-bound)」の世界です。後者は、直前のトークンから次の1トークンを紡ぎ出すために巨大な重み行列を一度だけ読み捨てる「メモリ帯域律速(Bandwidth-bound)」の世界です。同じ重みテンソルが、実行されるフェーズによって全く正反対の物理制約に晒されるというこの二重性こそが、静的な量子化フォーマットを機能不全に陥れる根本原因です。
3. 事件の登場人物
この物理と記号の衝突の最前線で、シリコンの限界に抗い、ビットの再編成を試みたエンジニアたちがいます。彼らはインテル(Intel Corporation)のHPC・AI高速化の心臓部であるパラレル・コンピューティング・ラボを拠点とし、長年にわたりコンパイラ、マイクロアーキテクチャ、線形代数カーネルの極限チューニングを共鳴させてきたチームです。
登場人物紹介
-
Evangelos Georganas(エヴァンゲロス・ジョルガナス)
英語表記:Evangelos Georganas / 現地語(ギリシャ語)表記:Ευάγγελος Γεωργάνας
推定年齢:39歳前後(2026年時点、存命中)。ギリシャ出身。カリフォルニア大学バークレー校(UC Berkeley)計算機科学専攻にてPh.D.を取得(指導教員:Katherine Yelick教授等、通信回避アルゴリズムの泰斗)。現在、Intel Labsのシニア・プリンシパル・エンジニア。世界最速の深層学習線形代数ライブラリ「LIBXSMM」のリードアーキテクトであり、CPUレジスタとキャッシュを分子レベルで制御するアセンブリコードの魔術師。本作においては、散乱する三元シンボルを1命令で整列させる「ビットの仕掛け人」として登場します。 -
Alexander Heinecke(アレクサンダー・ハイネッケ)
英語表記:Alexander Heinecke / 現地語(ドイツ語)表記:Alexander Heinecke
推定年齢:43歳前後(2026年時点、存命中)。ドイツ出身。ミュンヘン工科大学(TUM)にて情報科学のPh.D.を取得。ペタスケール地震シミュレーションでACMゴードン・ベル賞ファイナリストに名を連ねたHPCの巨頭。現在、Intel Fellow兼ディレクター。理論的な数理モデルを、XeonやIntel Xe GPUのシリコンダイ上で暴力的な実測スループットへと変換する推進役。本作では、三値モデルを冷徹なハードウェア物理へと押し出す冷徹なリアリストの役割を担います。 -
Pradeep Dubey(プラディープ・デュベイ)
英語表記:Pradeep Dubey / 現地語(ヒンディー語)表記:प्रदीप दुबे
推定年齢:64歳前後(2026年時点、存命中)。インド出身。パデュー大学にて計算機工学Ph.D.を取得。IEEE Fellow。数十年にわたりMMX、SSE、AVX、AMXといったIntelのSIMD/ベクトル/AI拡張命令のアーキテクチャ定義を主導してきた並列計算の生ける伝説。現在、Intel Fellow兼Parallel Computing Labディレクター。本作では、ゼロという沈黙の空間をアーキテクチャの武器へと転換する大局的な戦略家として描かれます。 -
敵対的査読者(The Hostile Reviewer)
本論文のプレプリント(arXiv:2609.16338)が学会へ投稿された際、査読者席の暗がりから容赦ない論理のメスを突き立てる仮想のPh.D.。彼らは「バッチ1のデコードだけで勝って何の意味があるのか」「自社の特殊命令に依存した局所最適ではないか」「最新のLunar Lakeで負けているではないか」と著者らを問い詰めます。彼らの疑義こそが、本研究を単なる「パッキング技法」から「PATE理論」へと飛翔させる最大の触媒となります。
4. 最初の謎
本書が解き明かすべき最初の謎は、極めて不条理な測定結果から始まります。
なぜ「より小さい」が「より速い」にならないのか:
5つの三元重みを1バイトに凝縮する高密度パッキング(5-trit方式)は、メモリフットプリントを極限まで削り取ります。しかし実測すると、重みをスカスカのまま2ビットで保持する素朴なフォーマットの方が、はるかに高いトークン生成レートを叩き出すことがあります。データを小さくしたはずの工夫が、なぜ実行パイプラインの首を絞めるのか。
なぜ速い機械ほど圧縮が不利になることがあるのか:
メモリ帯域が狭く貧弱な旧世代プロセッサでは劇的な高速化をもたらした圧縮カーネルが、最新鋭の広帯域メモリを積んだ超高速クライアントCPU(Lunar Lakeなど)に移植された途端、見る影もなく失速し、ベースラインに大敗を喫する。ハードウェアの進化がソフトウェアの最適化を無力化するという、アーキテクチャの裏切りはなぜ起きるのか。
なぜ同じ重みをCPUとGPUで違う姿にしなければならないのか:
サーバー用CPU(AVX-512)、クライアント用CPU(AVX2)、そしてGPU(Intel Xe2やNVIDIA Hopper)。同一の三元チェックポイントを実行しているにもかかわらず、なぜ最速のコードはそれぞれ全く相容れないアセンブリ命令列とメモリ配置を要求するのか。単一の「完璧なフォーマット」という幻想の崩壊が、ここに始まります。
5. 本書が最後に問い直すもの
これらの謎を巡る探求の旅は、単なる低レベルプログラミングのTips集では終わりません。我々が到達するのは、現代のAIエコシステム全体を規定している権力構造への根本的な疑問符です。
Quantization(量子化)ではなくRepresentation(表現): 静的に数値を丸めるだけの時代は終わりました。シリコンが要求するのは、動的に再構築可能な多次元の表現空間です。
Representation(表現)ではなくExecution(実行): 表現の優劣は、それ単体では決定されません。どのようなパイプラインで実行され、どこにレイテンシが隠蔽されるかという動的挙動がすべてを支配します。
Execution(実行)ではなく、誰が実行方式を決めるのか: 特定のGPUベンダーが提供する巨大で高価なプロプライエタリ環境に、世界中の知能インフラが平伏する現状は不可避なのか。三元適応表現という武器を手にしたとき、我々が日常的に手にする安価なコモディティチップが、巨大なAIデータセンターの軛(くびき)から推論の自由を奪還する道筋が見えてきます。
第一部 記号と物理の断絶――1.58ビットの神話とハードウェアの現実
第1章 1.58ビットはどこから来たのか――三つの記号と一つの神話
1.1 三元数という小さな宇宙
1.1.1 {-1, 0, +1}という世界
三値論理および三元重みテンソルが持つ本質的な魅力は、その極限的な禁欲性にあります。従来のディープラーニングにおいて支配的であった単精度浮動小数点(FP32: 32ビット)や半精度浮動小数点(FP16/BF16: 16ビット)は、仮数部と指数部を用いて無限に近い連続的な数値空間を表現しようと試みます。しかし、大規模言語モデルの莫大なパラメータ群が真に担っている役割が「特徴量の選択」と「シグナルの方向付け(抑制か促進か)」であるとするならば、重み行列の各要素に必要な情報は、極論すれば「通すな(0)」「強めよ(+1)」「反転せよ(-1)」の3状態に集約できるはずです。
この着想を大規模言語モデルの全結合層に適用したのがBitNetの系譜です。重み行列 W から乗算(Multiplication)という最もシリコン面積と電力を浪費する演算を完全に追放し、純粋な加算と減算、そしてゼロによる選択的スキップだけで推論を完結させる。この単純化は、チップの論理設計者にとってまさに夢のような跳躍でした。
1.1.2 log2(3)の意味
数学的に厳密な議論を進めましょう。アルファベット集合 Σ = {-1, 0, +1} の要素数が |Σ| = 3 であるとき、各シンボルが無作為かつ等しい確率(p = 1/3)で出現すると仮定するならば、クロード・シャノンが定式化した情報エントロピー H は次の式によって一意に決定されます。
H = - ∑i=13 pi log2 pi = - 3 × (1/3 log2 (1/3)) = log2 3 ≈ 1.5849625... [bits/symbol]
この log2(3) という数値は、いかなる可逆圧縮アルゴリズムを用いても、均一に出現する三つの記号列をシンボルあたり1.58496ビット未満の平均長で表現することは不可能であるという「情報理論的な下限(Theoretical Lower Bound)」を示しています。BitNet b1.58がその名称に「1.58」を掲げた背景には、自分たちの手法がこの理論的極限の際(きわ)に到達した究極の表現形式であるという強烈な自負が込められていました。
1.1.3 1.585という数字の誘惑
しかし、この1.585という数値は、同時に多くの研究者を思考停止に陥れる認知の罠(トラップ)となりました。論文誌やテックメディアには「1.58ビットで動くLLM」「重みメモリが16ビットから10分の1に激減する」という見出しが躍り、あたかもディスク上のファイルサイズやメモリバスのトラフィックが、そのまま元の1.58/16(約9.9%)に縮退するかのような幻想が振り撒かれました。人々は、シャノンの定理が前提としている「無限長のシンボル列に対する漸近的限界」と、「有限幅のバス幅を持つ物理回路でのトランザクション」の区別を忘れ去ってしまったのです。
1.2 「理論的に可能」と「実装可能」のあいだ
1.2.1 Entropyと固定長符号
情報理論を現実に引き戻す最初の冷や水は、可変長符号と固定長符号の断絶です。シンボルごとに fractional(端数)なビット幅を割り当てることは物理的に不可能です。単一の三元シンボルを格納しようとすれば、二進数の世界では最小の整数である2ビットを消費せざるを得ません。このとき消費されるビット幅は 2.000 bits/weight であり、log2(3) に対して実に 0.415ビット(約26.2%) の純然たる空間的無駄が生じます。
1.2.2 1.58-bitはinstruction formatではない
プロセッサの命令セットアーキテクチャ(ISA: Instruction Set Architecture)の歴史において、非整数ビット幅をネイティブに解釈するハードウェア命令など存在したことはありません。x86であれARMであれRISC-Vであれ、レジスタファイルの読み出しポートは8ビット、16ビット、32ビット、あるいはSIMDベクトルの128ビット、256ビット、512ビットといった単位で物理的に配線されています。1.58ビットという数値は、アルゴリズムの帳簿の上にのみ存在する記号論的な残余であり、ハードウェアの実行パイプラインが直接貪り食うことのできる「命令フォーマット」では断じてあり得ないのです。
1.2.3 Alignmentという現実
さらに過酷な現実がアライメント制約です。端数を解消するために複数のトリットを束ねるパッキング手法を考えましょう。代表的なものが「5-trit per byte」方式です。3の5乗は 243 であり、これは1バイト(8ビット = 256状態)に収まります。
35 = 243 ≤ 28 = 256 ⇒ 8 / 5 = 1.600 [bits/weight]
一見すると、これは1.585に肉薄するスマートな解決策に見えます。しかし、深層学習の行列乗算カーネルは、キャッシュ効率とSIMDベクトル化を極大化するために、128や256といった「2の冪乗(Power of Two)」のブロックサイズで重みを管理します。128要素の重みブロックを5トリットずつバイトに詰めようとすると、128は5で割り切れません。
⌈ 128 / 5 ⌉ = 26 [bytes] ⇒ 26 × 8 / 128 = 1.625 [bits/weight]
ブロック境界を跨ぐ複雑なビット跨ぎを排除し、アライメントを維持するためにパディングを挿入した結果、実効ビット幅は容赦なく 1.625 bits/weight へと跳ね上がります。実装という名の摩擦が、理論の輝きを削り落としていく最初の瞬間です。
1.3 三元LLMという発明
1.3.1 Ternary Weight
そもそも、なぜこれほどの実装上の摩擦を犯してまで、三元重みモデルを追求する価値があるのでしょうか。それは、推論時における浮動小数点FMA(Fused Multiply-Add: 積和演算)ユニットの完全なバイパスという、途方もない報酬が約束されているからです。
重みが {-1, 0, +1} の三値に制限されている場合、活性化ベクトル x と重み列 w の内積は、次のように展開されます。
y = ∑i xi × wi = ∑{i | wi = +1} xi - ∑{j | wj = -1} xj
乗算器は完全に姿を消し、残るのは単純なアキュムレータ(加算器)への入力のルーティングのみとなります。これにより、シリコンチップ上の演算器面積は激減し、演算に伴う動的消費電力(Dynamic Power Dissipation)は劇的に抑制されます。
1.3.2 QuantizationからWeight Distributionへ
しかし、近年の量子化対応学習(QAT: Quantization-Aware Training)や事後量子化(PTQ: Post-Training Quantization)の進化を詳細に追跡していくと、より深遠な事象がパラメータ空間で生起していることが明らかになってきました。重みを単に粗く丸める(Quantizeする)のではなく、ネットワーク全体の重み分布そのものが、学習プロセスを通じて三つの孤立したアトラクタへと収束していくダイナミクスです。
1.3.3 Zeroが意味するもの
そして、この三元分布の中で決定的な役割を果たしているのが「ゼロ(0)」です。従来の二値化ニューラルネットワーク(BNN: {-1, +1})が激しい精度崩壊(Accuracy Collapse)に苦しんできた最大の理由は、重みに「沈黙する権利」を与えなかった点にあります。特徴量を伝播させない、ノイズを遮断するというニューロンの本質的な抑制機能を果たすためには、ゼロという状態が絶対に必要なのです。後述するように、この「ゼロの存在」こそが、静的な情報理論の壁を粉砕する最大のトロイの木馬となります。
1.4 1.58-bit Barrierをどう再定義するか
1.4.1 Barrierはビット幅なのか
我々は問い直さねばなりません。「1.58ビットの壁」とは、本当にストレージメディア上に記録されるビット数の多寡のことなのでしょうか。ハードウェアにとって、メモリに1.58ビットで詰まっているか、1.625ビットで詰まっているかという差は、バス帯域の観点からは高々数パーセントの差異に過ぎません。
1.4.2 Storage BarrierとExecution Barrier
真の障壁は、ストレージ上のフットプリントを規定する「ストレージ・バリア(Storage Barrier)」ではなく、その詰まったビット列をレジスタへ引き上げ、演算器が咀嚼できる形式へと復元する際のレイテンシを規定する「エグゼキューション・バリア(Execution Barrier: 実行の壁)」です。解凍のための命令が実行パイプラインの実行ポートを埋め尽くし、メモリからデータが届いているにもかかわらず計算が進まない状態に陥るならば、ビット幅の微小な削減など何の恩恵ももたらしません。
1.4.3 本当の壁はどこにあるのか
本当の壁は、「記号の表現形式」と「実行装置の物理構造」が乖離していることそのものにあります。情報理論者が紙の上に描いたシンボルエントロピーと、シリコン設計者がリソグラフィで焼き付けたメモリ階層。この二つの世界の境界線こそが、三元LLM推論が打破すべき真のフロンティアなのです。
コラム:筆者の現場録――秋葉原の片隅で考えた「失われた0.415ビット」
かつて筆者が初期の組み込み向けSIMDカーネルを夜を徹してデバッグしていた頃、アセンブラの吐き出す命令ストール(Stall)に頭を抱えていた。情報理論の教科書を読めば、ハフマン符号や算術符号を使えばデータは極限まで縮むと書いてある。意気揚々と可変長デコード処理をC言語で書き、インラインアセンブラでねじ込んだ。結果はどうだったか。素朴に16ビット幅で並べたコードに対して、実行速度は3分の1に叩き落とされたのだ。分岐予測は外れまくり、ビットシフト命令がALUの単一ポートを奪い合っていた。チップのダイ写真を顕微鏡で覗き込んだとき、痛感した。シリコンの神様は、美しいエントロピーなど求めていない。整然と並んだ、頭の悪いバイト列を猛烈な勢いで欲しがっているのだ、と。
第2章 数字を読む機械――CPUは「1.58ビット」を知らない
2.1 メモリからレジスタまで
2.1.1 Cache Line
現代の計算機が主記憶装置(DRAM)と通信する際の基本通貨は、決して単一のバイトではありません。標準的なx86プロセッサやハイエンドARMプロセッサにおける最小トランザクション単位は、連続する64バイト(512ビット)のキャッシュライン(Cache Line)です。CPUコアがメモリ上のわずか1ビットのフラグを読み出そうとするだけでも、メモリコントローラは周辺の64バイト全体をごっそりとL3キャッシュ、L2キャッシュ、そしてL1キャッシュへと吸い上げます。重みがどれほど巧妙に1.58ビットにパッキングされていようと、その読み出し要求がキャッシュラインの境界を跨ぐ(Misaligned Access)瞬間、プロセッサは2倍のメモリアクセス要求を発行し、性能の断崖絶壁へと転落します。
2.1.2 SIMD Register
キャッシュから引き上げられたデータは、プロセッサの深部にあるレジスタファイルへと送られます。近年のAI推論を牽引するSIMD(Single Instruction, Multiple Data: 単一命令複数データ処理)レジスタ――例えばIntel AVX-512における512ビット幅の zmm レジスタ群――は、内部が極めて厳密なレーン構造(32ビット整数なら16レーン、16ビット浮動小数点なら32レーン、8ビット整数なら64レーン)に分割されています。プロセッサがベクトル演算器を全開で駆動させるためには、レジスタの各レーンに整然と数値が配置されていなければなりません。奇妙なビット幅のデータは、演算器の手前で、この固定長レーンへと「荷解き(アンパック)」される運命にあります。
2.1.3 Load/Decode/Compute
パイプラインの観点から見れば、低ビット推論カーネルの内部ループは常に三者の綱引きです。
- Load: 圧縮された重みバイト列をメモリ/キャッシュからSIMDレジスタへロードする。
- Decode (Unpack): ビット操作命令を駆使して、圧縮シンボルを演算可能な固定長形式へと復元する。
- Compute: 復元された重みと活性化ベクトルの間で内積を計算する。
これら三つのステージのうち、どれか一つでも他を圧倒して時間を消費すれば、残りの二者は遊休状態(Idle)に追い込まれます。
2.2 Memory Wallの再来
2.2.1 Arithmetic Intensity
この力学を定量化するための絶対的な指標が、算術強度(Arithmetic Intensity: AI)です。算術強度は、メモリシステムから1バイトのデータを転送する間に、演算器が何回の浮動小数点演算(あるいは整数積和演算)を実行できるかという比率として定義されます。
Arithmetic Intensity (AI) = Useful Operations [FLOPs or OPs] / Data Transferred [Bytes]
2.2.2 Roofline Model
サミュエル・ウィリアムズ(Samuel Williams)らが2009年に提唱したルーフラインモデル(Roofline Model)は、プロセッサの達成可能性能 Pmax を、ハードウェアの理論最大演算スループット Π [FLOP/s] と、持続可能メモリ帯域幅 β [Bytes/s]、そしてワークロードの算術強度 AI の関数として美しく図式化します。
Pmax = min ( Π, β × AI )
AIがプロセッサの変曲点(Ridge Point: Π / β)よりも低い領域では、演算器がどれほど余っていようとも、性能はメモリ帯域幅によって完全に頭打ちになります(メモリ帯域律速)。逆にAIが変曲点を超えると、性能は演算器のピーク性能によって制限されます(計算律速)。
2.2.3 Bandwidth versus Throughput
現代のハイエンドGPUやサーバー向けCPUにおいて、この変曲点は極めて高い位置にあります。例えば、1秒間に数テラフロップスからペタフロップスを叩き出す演算器に対し、提供されるメモリ帯域幅は数百GB/sから数TB/sに過ぎません。結果として、変曲点は数十から数百 FLOPs/Byte という極めて高い値に設定されています。
2.3 DecodeとPrefillの二つの時間
2.3.1 Batch=1の世界
ここで、LLMの自己回帰デコード(Autoregressive Decode)を考えます。単一のユーザーがチャットボットと対話している状況(バッチサイズ N = 1)では、1つの新規トークンを生成するために、モデルのすべての重み行列 W ∈ RK × M が1回だけメモリから読み出され、1つの入力ベクトル x ∈ R1 × K と掛け合わされます。
このとき実行される演算量は 2 × K × M FLOPs であり、読み出される重みデータ量は(FP16ならば)2 × K × M バイトです。したがって、算術強度はわずか 1.0 FLOP/Byte にしかなりません。変曲点が100を越える現代のプロセッサにおいて、AI = 1.0 という値は、演算器の99%がメモリからのデータ到着を待って虚無の時間を過ごしていることを意味します。デコード相は、徹底的かつ絶望的なまでにメモリ帯域律速(Bandwidth-bound)なのです。
2.3.2 Batchが増えると何が変わるのか
しかし、サーバーが複数のユーザーからの要求を束ね、バッチサイズを N = 32 や N = 64 へと拡大した瞬間、世界の物理法則は一変します。メモリから1度読み出された同一の重み行列 W が、N個の異なる入力ベクトルに対して使い回される(Reuseされる)ため、算術強度はバッチサイズに比例して N 倍へと跳ね上がります。
AI(N) ≈ (2 × N × K × M) / (bw × K × M) = 2N / bw
Nが十分に大きくなれば、AIはプロセッサの変曲点を軽々と飛び越え、処理は計算律速(Compute-bound)の領域へと突入します。ここでは、もはやメモリ帯域の節約は何の価値も持ちません。演算器をいかに淀みなく回転させ続けるかだけが支配的な関心事となります。
2.3.3 Weight Reuseという第三の変数
同様の現象は、長いコンテキストを一度に読み込む「Prefill相」でも発生します。プロンプトのトークン長が L = 2048 であれば、一度のGEMM(一般行列乗算)によって重みは2048回再利用されます。重みの再利用係数 R(Weight Reuse Factor)こそが、ハードウェアのボトルネックをメモリ側から演算器側へと強制的にスライドさせる決定的なレバーなのです。
2.4 同じモデル、違う物理
2.4.1 Server CPU
Intel Xeon Platinum 8592+(Emerald Rapids: EMR)のような64コアのサーバー向けCPUは、高密度のDDR5メモリチャネル(8チャネル)を備え、理論帯域幅は約300 GB/sに達します。しかし、64もの強力なコアが同時に群がれば、コアあたりに割り当てられる帯域幅はわずか 4.7 GB/s 程度に希釈されます。コア内部のAVX-512 FMA演算能力に対して帯域が圧倒的に不足しているため、サーバーCPUはバッチ1の推論において強烈なメモリ帯域飢餓状態に置かれます。
2.4.2 Client CPU
一方、Intel Core Ultra 7 258V(Lunar Lake: LNL)のような最新クライアント向けSoCは、極めて特異な物理バランスを提示します。わずか8個のCPUコア(4つの高性能Pコアと4つの高効率Eコア)に対し、オンパッケージの超高速LPDDR5Xメモリが直結され、実測で100 GB/sを超えるストリーミングリード帯域を叩き出します。コアあたりに換算した利用可能帯域幅は、サーバーCPUの数倍に達する 12〜13 GB/s/core です。ここでは、メモリ帯域はもはやボトルネックではありません。むしろ、少ないコアがその広大な帯域から雪崩のように押し寄せるデータをアンパックしきれるかという、命令処理能力の上限が露わになります。
2.4.3 GPU
GPU(NVIDIA H100やIntel Arc Pro B70など)は、テラバイト級のHBMや広帯域GDDR6メモリによって武装していますが、数千から数万に及ぶ超並列スレッドが同時に走るため、個々のスレッドから見たレジスタ容量やローカルメモリ(Shared Memory)の帯域幅は極めてタイトに制約されます。
2.4.4 Edge Accelerator
さらに車載や産業機器に組み込まれるエッジAIアクセラレータでは、消費電力が数ワットから十数ワットに制限され、メモリバス幅は64ビットや32ビットに切り詰められています。物理プラットフォームごとに「何が余っていて、何が足りないのか」というバランスシートは、完全に別物なのです。
コラム:プロファイラの囁き――キャッシュミスと結婚した男
Linuxの perf コマンドを叩いて、PAGE-FAULTS や LLC-LOAD-MISSES のカウンターが猛烈な勢いでカウントアップしていくのを眺めるのが好きだ、と言うと大抵の同僚は気味の悪いものを見る目で筆者を見た。しかし、ハードウェアの真実はいつだって性能カウンターの冷え切った数字の中にしかない。ある日、最適化したはずの低ビットGEMVカーネルのIPC(Instructions Per Cycle)が0.4を下回っていた。1サイクルに4命令実行できるはずの最新コアが、10サイクルに4回しか息をしていない。犯人はメモリコントローラのキュー詰まりだった。数字を読む機械は、プログラマの意図など1ミリも忖度しない。与えられたバイト列を、決められた配管の太さでしか通せないのだ。
第3章 敵対的査読者の反乱――「それはDecodeでしか速くない」
3.1 第一の異議――Batch=1問題
3.1.1 Decode専用ベンチマークではないのか
さて、ここで学術界の暗闇から現れる「敵対的査読者」の厳しい詰問に耳を傾けなければなりません。彼らは、新しい量子化レイアウトを提案する論文のプレプリントを眺め、即座にその急所を見抜きます。
「著者の主張するスループット向上は、バッチサイズ1の自己回帰デコードという、極端に偏った条件下でのみ測定された砂上の楼閣ではないか?」
3.1.2 Prefillでunpackingが支配する可能性
彼らの論理は冷徹です。LLMの実用的なサービング環境において、システムのスループットを真に決定づけるのは、入力プロンプトを高速に消化するPrefillフェーズ、および多数の同時リクエストを束ねる連続バッチ処理(Continuous Batching)です。これらの領域では、前述の通り算術強度 AI が急激に上昇します。
メモリ帯域が律速でなくなった世界において、重みを解凍するために挿入された無数のビット操作命令(シフト、マスク、テーブル参照)は、純粋なオーバーヘッドとしてパイプラインのクリティカルパスに重くのしかかります。「デコードを20%高速化するために、Prefillのレイテンシを2倍に悪化させているのではないか?」という指摘は、実用システムを見据える技術者にとって逃れることのできない致命的な問いです。
3.1.3 Batch inferenceという別世界
商用のクラウド推論基盤において、GPUクラスタはバッチサイズを32、64、あるいは128へとスケールさせて稼働率を限界まで高めます。バッチ1という「エッジ端末の片隅でのみ許された特殊なユースケース」に過剰適合したアルゴリズムは、ハイパースケーラーのデータセンターでは産業廃棄物に等しいのではないか。この痛烈な反論に対し、従来の低ビット提案の多くは口を噤んできました。
3.2 第二の異議――Lunar Lake問題
3.2.1 帯域が余っている機械
第二の異議は、さらに具体的かつ残酷です。Intel自身の最新チップである「Core Ultra 7 258V(Lunar Lake)」での実測結果という、逃れようのないファクトに基づく批判です。
第2章で触れた通り、Lunar Lakeはクライアント向けSoCでありながら、8個のコアに対して100 GB/sを超える圧倒的なメモリ帯域幅を誇ります。メモリからデータを運ぶパイプが太すぎるため、バッチサイズ1のデコード時であっても、コアはメモリ待ちで苦しむことがありません。
3.2.2 少コア・広帯域という逆風
このプラットフォーム上で、著者らの提案する高圧縮レイアウト(BITCOS)を実行すると何が起きたか。原論文のTable IIおよびFigure 11(c)が正直に告白している通り、BITCOSカーネルは、単純な2ビット固定長カーネルに対して完全な敗北を喫しました。
コア数が少ないため、解凍命令を並列に処理する絶対的な演算能力が不足し、メモリ帯域を使い切る前にコアの命令デコーダと実行ポートがパンクしたのです。メモリからデータはすでに届いているのに、ビットをほどく作業が追いつかず、パイプラインが停止する。小さく圧縮したはずのデータ形式が、広帯域・少コアという新しいハードウェアトレンドの前に無惨にも討ち死にした瞬間でした。
3.2.3 Compressionが逆効果になる瞬間
敵対的査読者は嘲笑します。「君たちのアルゴリズムは、ハードウェアが進歩してメモリ帯域が広くなればなるほど遅くなるのか?未来のアーキテクチャに逆行する最適化に、一体何の学術的価値があるというのかね?」と。
3.3 第三の異議――PDEP問題
3.3.1 「Intelで動く」は「一般に動く」ではない
第三の異議は、可搬性(Portability)に対する構造的疑義です。著者らがx86 CPU上でBITCOSの超高速解凍を実現した中核兵器は、BMI2(Bit Manipulation Instruction Set 2)に含まれる PDEP(Parallel Bits Deposit: 並列ビットデポジット)命令でした。
しかし、査読者は知っています。PDEP は長年、Intelプロセッサ特有の飛び道具であり、ライバルであるAMDのZen 1やZen 2マイクロアーキテクチャにおいては、専用のハードウェア回路を持たず、マイクロコードによる数百サイクルのエミュレーションとして実装されていた悪名高い命令です。Intelの最新Xeon上でどれほど華々しい数字を叩き出そうとも、世界中のサーバーの半分を占めるAMD環境や、世界中のスマートフォンとエッジ端末の99%を支配するARMプロセッサ上で動かなければ、それは普遍的な技術とは呼べません。
3.3.2 Xe2 SLM依存という疑惑
同様に、GPU実装において著者らが提示したXe2アーキテクチャ向けのカーネルも、共有ローカルメモリ(SLM)内の2KBルックアップテーブルという、極めてプラットフォーム依存の強い構造に頼っています。NVIDIAのTensor CoreやAMDのCDNAアーキテクチャにおいて、同様のメカニズムが成立するという保証はどこにあるのか。
3.3.3 ARMとNVIDIAはどうするのか
「ISA(命令セット)の局所的な珍しい命令に魂を売り渡したハックを、普遍的なブレークスルーと呼ぶのは誇大広告ではないか?」――この査読者の銃口は、ソフトウェア最適化が常に抱える「汎用性と性能のトレードオフ」という急所を正確に捉えていました。
3.4 査読者が見落としている問い
3.4.1 「BITCOSは速いか」ではない
これらの批判は極めて正当であり、生半可な言い訳では突破できません。しかし、査読者たちもまた、一つの巨大なパラダイムの檻の中に閉じ込められていました。彼らは依然として、「BITCOSという単一のフォーマットが、あらゆる条件下でベースラインより速いか否か」という二元論の土俵の上で評価を下そうとしていたのです。
3.4.2 「いつBITCOSを使うべきか」
問うべき真の問いは、「BITCOSは速いか」ではありません。
「いかなるハードウェアバランスとワークロード条件下において、BITCOSは数学的必然として勝利し、そしていかなる条件下で他の表現にバトンを渡すべきなのか?」
3.4.3 そして「いつ使わないべきか」
敗北を隠蔽するのではなく、敗北する境界線を数理的に厳密に予測し、システムが自動的に最適な別の表現へと切り替える自律機構を構築すること。査読者の反乱は、固定的な「フォーマット至上主義」を葬り去り、動的な「相適応型アーキテクチャ(PATE)」へと至る扉を押し開く決定的な契機となったのです。
コラム:査読用コメント欄の向こう側――「Reject」の文字を見た夜
国際学会の通知メールを開き、Reviewer 2の欄にぎっしりと書き込まれた冷酷なパラグラフをスクロールするときの胃の収縮は、何度経験しても慣れるものではない。「The proposed format lacks general applicability and degrades significantly under compute-bound regimes...」彼らの言う通りだった。図星だからこそ、血圧が上がる。深夜のオフィスでホワイトボードの前に立ち、消しゴムを叩きつけそうになりながら考えた。待てよ、と。彼らは「万能ではない」ことを突いている。ならば、「万能のフォーマットなど存在しないこと」を数学の定理として証明してしまえば、彼らの批判はすべてこちらの理論の正しさを補強する実験データに化けるのではないか? 敵を味方に変える論理の反転は、いつだって最悪の査読レポートから生まれる。
第4章 歴史はすでに知っていた――最適化はいつも反転する
4.1 CISC vs RISC
4.1.1 命令を複雑にする
計算機科学の歴史を振り返れば、表現と実行の綱引きによって最適解が真逆に反転するというドラマは、何十年も前から繰り返されてきた古典劇に過ぎません。1970年代後半のCISC(Complex Instruction Set Computer: 複合命令セット計算機)の隆盛がその筆頭です。
当時、主記憶装置(磁気コアメモリや初期の半導体RAM)は天文学的に高価であり、極めて低速でした。したがって、システムの至上命題は「プログラムのメモリ占有量を1バイトでも削り、1回の命令フェッチで可能な限り多くの仕事をさせること」にありました。VAX-11に代表されるCISCプロセッサは、多項式の評価や文字列の走査を単一の複雑怪奇な機械語命令で実行できるように設計されました。コード密度(Code Density)こそが正義であり、メモリフットプリントの最小化がそのまま性能の最大化を意味していた時代です。
4.1.2 命令を単純にする
しかし、半導体リソグラフィ技術の爆発的な進歩によってDRAMの集積度が向上し、プロセッサ内にオンチップの命令キャッシュが搭載され始めた1980年代、デビッド・パターソン(David Patterson)やジョン・ヘネシー(John Hennessy)らが率いるRISC(Reduced Instruction Set Computer: 縮小命令セット計算機)革命が勃発します。
彼らが看破したのは、複雑な命令をデコードするために浪費される制御論理回路のオーバーヘッドでした。命令フォーマットを固定長の単純なものに統一し、デコードを極限まで簡素化してパイプラインを高速に回せば、プログラムの総バイト数が多少増えようとも、実際の実行時間は圧倒的に短縮される。命令密度の最大化から、実行パイプラインの平滑化へ。物理の重心が移動した瞬間、最適化の絶対的正義は180度反転したのです。
4.1.3 「どちらが正しいか」という間違った問い
後世の我々が知る通り、「CISCとRISCのどちらが普遍的に正しかったのか」という問い自体が無意味でした。現代のx86プロセッサは、外側は可変長のCISC命令セットを装いながら、内部のハードウェアデコーダがそれを単純な固定長の内部命令(μOPs: マイクロオペレーションズ)へとオンザフライで分解して実行するハイブリッド構造を採用しています。静的な命令フォーマットの純血主義など、シリコンの物理の前には雲散霧消したのです。
4.2 VLIWの挫折
4.2.1 コンパイラに未来を予測させる
同様の痛切な教訓は、1990年代末から2000年代初頭にかけてIntelとHPが社運を賭けて開発したIA-64(Itanium)アーキテクチャ、すなわちVLIW(Very Long Instruction Word)の興亡からも読み解くことができます。
VLIWの思想は、ハードウェアの実行時スケジューリング論理を全廃し、コンパイル時に高度な静的解析を行って、依存関係のない複数の命令を巨大な命令語(エクスプリシット・パラレル・インストラクション・コンピューティング: EPIC)へとパックすることにありました。「実行時にハードウェアで並列性を抽出するのは重すぎる。静的にすべてを決めておけば、実行器は極限までシンプルかつ高速になる」という信仰です。
4.2.2 ハードウェアが変わった瞬間
しかし、この壮大な構想は現実の壁に激突して脆くも崩壊しました。コンパイラがどれほど知能を尽くして静的なスケジュールを組もうとも、実行時の動的な物理事象――キャッシュミスによる予期せぬレイテンシ、メモリアクセスの競合、入出力割り込み――をコンパイル時に完全に予測することは不可能だったからです。一度キャッシュミスが発生すれば、後続の美しく整列された命令群はすべて道連れとなってパイプライン上で凍りつきました。
4.2.3 静的最適化の限界
ハードウェアの世代が変わり、キャッシュレイテンシやパイプライン段数がわずかに変化しただけで、過去のコンパイル結果は最適性を失い、再コンパイルなしには性能が出ない。VLIWの挫折が歴史に刻んだ教訓は明白です。「実行時の動態を無視して、静的に決定されたフォーマットに過剰適合したシステムは、ハードウェアの進化と動的環境の変化に耐えられない」という真理です。
4.3 GPU memory wall
4.3.1 FLOPS競争
近年のグラフィックスおよびディープラーニングアクセラレータの軌跡もまた、この反転の繰り返しでした。2010年代、GPUは純粋な浮動小数点演算器(ALU)をシリコンダイの許す限り敷き詰め、テラフロップス競争に明け暮れました。
4.3.2 Bandwidth競争
しかし、モデルサイズがパラメータ数十億の壁を突破したとき、演算器はどれほど高速化されてもデータを供給されない飢餓状態に陥りました。性能競争の主戦場は、FLOPSからメモリ帯域幅(GB/s)へと不可逆的にシフトしたのです。
4.3.3 HBMという救済策
この帯域飢餓に対するシリコンバレーの回答が、シリコン貫通電極(TSV: Through-Silicon Via)を用いてDRAMダイを三次元に積層する超高価なソリューション、HBM(High Bandwidth Memory)でした。帯域幅は飛躍的に拡大しましたが、それはパッケージングコスト、歩留まり、消費電力、そしてサプライチェーンの極端な脆弱性という、新たな絶望的な壁を出現させることになりました。
4.4 暗号通信の最適化史
4.4.1 ビットを詰める
通信とセキュリティの領域におけるデータ表現の変遷も、示唆に富んでいます。初期の暗号通信プロトコルでは、ネットワーク帯域が貴重であったため、パケットヘッダやペイロードはビットフィールド単位で細密にパッキングされていました。
4.4.2 ビットを展開する
しかし、ギガビットイーサネットが普及し、パケット処理のボトルネックが伝送路の帯域からCPUのパケット解析処理能力へと移るにつれ、ビットを詰める処理そのものがルータの処理限界を規定するようになりました。プロトコルはアライメントを重視したバイト指向、ワード指向へと回帰していきました。
4.4.3 AES-NIが変えたもの
そして決定的な転換点は、Intelが暗号処理のための専用命令セット「AES-NI」をシリコンに刻んだ瞬間です。かつてソフトウェアで数千サイクルを要していた複雑なS-Box置換とビット攪拌が、単一の命令で数サイクルで完了するようになった途端、暗号化処理のコストはゼロ同然に吹き飛びました。ハードウェア命令の追加が、ソフトウェアのアルゴリズム選定の前提条件を一夜にして無価値にしたのです。
4.5 歴史から得られる一つの原則
4.5.1 表現は命令セットに従属する
これらの歴史的事例が異口同音に指し示している原則は、ただ一つしかありません。
「最適なデータ表現とは、データの持つ数学的性質だけで決まるものではない。それは、そのデータを読み解く命令セットの構造に絶対的に従属する。」
4.5.2 命令セットは物理に従属する
そして、その命令セットの有効性を規定するのは、シリコンダイ上のトランジスタ予算、配線遅延、メモリコントローラのバス幅、キャッシュ階層の容量といった抗いようのない物理構造です。
4.5.3 物理が変われば最適表現も変わる
物理が変われば、命令のコストが変わり、命令のコストが変われば、最適な表現は反転する。三元LLMの重みをどのようにパッキングすべきかという問題も、この歴史の鉄則から逃れることはできません。固定的な「1.58ビット」という静的フォーマットの呪縛を解き放ち、物理の変動に合わせて姿を変える柔軟な表現体系を構築しなければならない所以が、ここにあります。
コラム:博物館のシリコン――VAXとItaniumの眠る場所
カリフォルニア州マウンテンビューのコンピュータ歴史博物館の静まり返った展示室を歩くのが好きだ。そこには、かつて世界最高の頭脳たちが「これこそが計算機の究極の未来だ」と確信して設計したプロセッサたちが、ガラスケースの中で静かに眠っている。数千個のマイクロコードで埋め尽くされたVAXの基板。巨大な命令スロットを誇らしげに掲げたItaniumのダイ。彼らが間違っていたのではない。彼らの設計思想は、当時の、あるいは彼らが夢想した物理環境においてはこの上なく正しかったのだ。ただ、物理が彼らを追い越していった。いま我々が書いている低ビットのカーネルコードも、10年後のエンジニアから見れば「なぜ彼らはこんな奇妙なビット操作に必死になっていたのか」と微笑まれる遺物に過ぎないのかもしれない。だがその悪戦苦闘の軌跡だけが、次の物理の扉を開けるのだ。
第二部 ゼロの幾何学――BITCOSレイアウトと命令レベルの錬金術
第5章 ゼロは「何もない」のではない――三元重みの地形を測る
5.1 Zero Density
5.1.1 zという新しい座標
第一部において我々は、三元シンボル {-1, 0, +1} を等確率として扱う情報理論的アプローチの限界を論じました。では、現実の三元大規模言語モデルの内部では、一体どのような現象が起きているのでしょうか。
Intelの研究チーム(Georganasら)は、推論カーネルの設計に入る前に、極めて地道かつ決定的な実証的調査を行いました。公開されている最先端(SOTA: State-of-the-Art)の三元LLM、計7つのモデルファミリー、実に29種類ものチェックポイントを収集し、その全重みテンソルに含まれるシンボルの出現頻度を徹底的にカウントしたのです。
調査対象には、スクラッチから4兆トークンで事前学習された三元モデルの基準点「BitNet b1.58 2B4T」、1.7Bから27Bに至る高密度モデル群「Bonsai」、Qwen3をベースとした事後量子化モデル「CAT-Q」(MoEモデルを含む)、低ビットQATの研究から生まれた「ParetoQ」、99Mから3.9Bまでを網羅する「TriLM (Spectra)」、さらにはMoE推論モデル「Maple 20B-A1B」、そして「BitCPM-CANN」が含まれていました。
そこで得られた測定結果は、従来の固定長パッキングの前提を根底から覆すものでした。重みテンソルを構成する全要素の中で、「ゼロ(0)」が占める割合(ゼロ密度: Zero Density z)は、最大で実に 51.5%(CAT-Q Qwen3-1.7B)に達していたのです。
5.1.2 Dense TernaryとSparse Ternary
表Iに示された全29モデルのゼロ密度 z を俯瞰すると、驚くべき統計的規則性が浮かび上がります。最もゼロ密度が低いモデル(Bonsai 27B)であっても z = 29.66% を記録しており、全モデルの平均値はおよそ 40%前後 に収束していました。
三元モデルの重みは、決して {-1, 0, +1} の間を均等に揺れ動く「高密度な三値(Dense Ternary)」などではありませんでした。それは、全要素の4割から5割以上が完全に沈黙している、高度に疎な三値(Sparse Ternary)の地形を形成していたのです。
5.1.3 モデルごとに違う「ゼロの地図」
このゼロ密度 z は、モデルの学習手法によっても固有の指紋を示します。事後量子化(PTQ)によって構築されたCAT-Qファミリーは総じて高いゼロ密度(46%〜51%)を示す傾向があり、一方で大規模な事前学習を経たBonsaiやBitCPMは37%〜40%近傍に密集しています。重み空間の地形は均一ではなく、モデルの出自とアーキテクチャに応じて多様な「ゼロの山脈と谷」を描き出しているのです。
5.2 2-z bits/weight
5.2.1 Presence Bitmap
この「全要素の4割以上がゼロである」という統計的事実を前にしたとき、情報理論の風景は劇的に塗り替わります。ゼロの出現確率 p0 が 0.4 や 0.5 に達しているならば、もはや各シンボルを等確率として log2(3) ≈ 1.585 を参照点にする正当性は完全に失われます。
Georganasらは、この非対称性を極限まで単純なデータ構造へと変換するアイデアを着想しました。それが「BITCOS (BITmap and COmpacted Signs: ビットマップとコンパクト符号)」レイアウトです。テンソルを二つの独立したコンポーネントへと分解します。
- 存在ビットマップ(Presence Bitmap): 重みテンソルと全く同じ要素数を持つ1ビットのフラグ列。重み要素が非ゼロ(-1または+1)であれば「1」を立て、ゼロであれば「0」とする。
- コンパクト符号ベクトル(Compacted Sign Vector): 非ゼロの重み要素に対してのみ、その符号(正なら0、負なら1)を1ビットで記録し、テンソル順に隙間なくパッキングしたビット列。
5.2.2 Sign Stream
このレイアウトの真骨頂は、ゼロという要素が「符号ベクトルにおいて一切の空間を消費しない」という点にあります。全要素数 Nw に対し、存在ビットマップは正確に Nw ビットを消費します。そして、非ゼロ要素の総数は (1 - z) × Nw 個ですから、符号ストリームが消費するビット数は (1 - z) × Nw ビットです。
5.2.3 なぜゼロが圧縮率を決めるのか
したがって、重み要素1個あたりに割り当てられる平均ビット幅 B(z) は、次のような極めて優美な一次関数として定式化されます。
B(z) = 1 [bitmap] + (1 - z) [sign] = 2 - z [bits/weight]
何という単純明快さでしょう。ゼロ密度 z が上昇すればするほど、ストレージコストは線形に低下します。
- もしゼロが半分(z = 0.50)を占めるモデルであれば、実効ビット幅は 1.500 bits/weight となり、あの神話的な1.585ビットの壁をあっさりと下回ります。
- 最疎モデルであるCAT-Q Qwen3-1.7B(z = 0.5148)においては、純粋なシンボルレートは実に 1.485 bits/weight に到達します。
従来の5-tritパッキングの実効レートである 1.625 bits/weight と比較してみましょう。
2 - z < 1.625 ⇔ z > 0.375 (37.5%)
損益分岐点(Break-even Point)は、ゼロ密度 37.5% の位置にあります。実測された29モデルのうち、実に26モデルにおいて z > 0.375 が成立していました。すなわち、BITCOSレイアウトを採用するだけで、現存する三元LLMのほぼ9割が、従来のいかなるパッキング手法よりもコンパクトにメモリ上に格納できることが数学的に証明されたのです。
5.3 ゼロの空間構造
5.3.1 Random zero
しかし、データ構造が優れていることと、それが実機上で高速に解凍できることは同義ではありません。ハードウェアの観点から次に問われるべきは、これらゼロ要素が空間的にどのように分布しているかという「幾何学」です。
仮にゼロが完全にランダム(ベルヌーイ過程)に散らばっているとすれば、メモリ上の各ブロックに含まれる非ゼロ要素の数は激しく変動し、固定長の命令シーケンスで規則的に処理することは困難になります。
5.3.2 Clustered zero
ニューラルネットワークのパラメータを詳細に観察すると、ゼロは完全にランダムに分布しているわけではありません。特定の出力チャネルや特定の特徴量マップの周囲に、ある程度の塊(クラスター)を形成して局所化する傾向が見られます。
5.3.3 Layer-dependent zero
さらに、モデルの深さ方向においてもゼロ密度はダイナミックに変化します。入力に近い浅い層(Embedding直後の初期トランスフォーマーブロック)では、入力特徴量の多様性を維持するためにゼロ密度は比較的低く抑えられます(z ≈ 0.35〜0.40)。
一方、深い層、特にFFN(Feed-Forward Network)の中間層やDown-projection層においては、高度に抽象化されたスパースな表現が支配的となり、ゼロ密度は容易に 55%〜65% を超える高密度な沈黙の領域へと突入します。
5.3.4 Block-dependent zero
この層間・ブロック間の不均一性は、固定的な圧縮スキームにとっては頭痛の種ですが、後述する適応型の推論システムにとっては、局所的な最適化を仕掛けるための絶好のキャンバスとなります。
5.4 「ゼロ率」から「ゼロの幾何学」へ
5.4.1 同じzでも速さが違う
マイクロアーキテクチャの視座に立ったとき、我々は単なるスカラー値としての「ゼロ率」を超えて、「ゼロの幾何学(Information Geometry of Zeros)」を把握しなければなりません。なぜなら、全く同一のゼロ密度 z = 0.40 を持つ二つの重み行列であっても、その配置構造の違いによって、解凍カーネルの実行速度は最大で20%以上も変動し得るからです。
5.4.2 Locality
非ゼロ要素が32要素のベクトル幅の中に綺麗に収まっているか、それとも境界を跨いで激しく跨いでいるかという空間的局所性(Spatial Locality)。
5.4.3 Rank
ビットマップの任意の位置までにいくつの「1」が存在するかという階数(Rank)の計算が、キャッシュラインの境界内で完結するか否か。
5.4.4 Entropy
ゼロが織りなす空間の疎密パターンそのものが、プロセッサの分岐予測器やメモリアライナに対する物理的な負荷を決定づけます。ゼロとは、単なる「値が0であること」ではありません。それは、ハードウェアの実行リソースを再配分するための、極めて雄弁な構造的シグナルなのです。
コラム:統計の罠――平均値に殺されたプログラマ
「全層の平均ゼロ密度は42%です」。データサイエンティストから渡されたレポートのその一行を信じて、均一なタイルサイズで最適化したカーネルを書いたことがある。見事にクラッシュした。ある層ではゼロが20%しかなく符号バッファがオーバーフローし、別の層ではゼロが70%に達してALUが完全に飢餓状態に陥っていた。「平均水深1メートルの川で溺れ死ぬ」という有名な警句があるが、ニューラルネットワークのテンソル空間ほど、平均値が現実のデッドロックを覆い隠してしまう場所はない。現場のシステムプログラマは、いつだって最悪の局所分散に備えてコードを書かなければならないのだ。
第6章 BITCOS――二つのビット列で三つの世界を表す
6.1 PresenceとSign
6.1.1 Bitmap
BITCOSのアーキテクチャを、より解像度を高めて解剖しましょう。図2に示された 8 × 8 の重みテンソルの小宇宙を想像してください(全64要素)。
ビットマップは、行列の各列(あるいは各行)を自然な二進ワードとして表現します。例えば、ある列の8要素が (+1, +1, +1, +1, -1, 0, -1, +1) であったとすれば、ゼロである第5行(インデックスは0始まり)のみがビット0となり、残りの7要素がビット1となります。
Row Index: 0 1 2 3 4 5 6 7
Value: + + + + - 0 - +
Bitmap: 1 1 1 1 1 0 1 1 ⇒ 0xDF (16進数)
ビットマップの重み1ワード(32ビット整数)をロードするだけで、デコーダは一挙に32行分の「存在/非存在」の全貌を、単一のメモリアクセスで把握することができます。
6.1.2 Compacted Sign
一方、符号ベクトルは、非ゼロの要素に対応する符号ビットのみを、出現順に物理メモリ上に隙間なく詰め込みます。上記の例では、非ゼロ要素は7個存在し、それぞれの符号は「正(0), 正(0), 正(0), 正(0), 負(1), 負(1), 正(0)」です。したがって、符号ストリームには次の7ビットのみが記録されます。
Compacted Signs: 0 0 0 0 1 1 0 (7 bits)
不在である第5行のための符号ビットは、メモリ上に1ミリも存在しません。完全に消去されているのです。
6.1.3 Rank-select
このとき、任意の行 i の符号ビットを取り出すためには、古典的な簡潔データ構造(Succinct Data Structure)の基本操作である Rank操作 が必要となります。行 i より手前に、一体いくつの非ゼロ要素が存在していたか。
ri = Rank1(Bitmap, i) = popcount( Bitmap & ((1 << i) - 1) )
この Rank ri こそが、符号ストリーム内におけるその要素の正確なビットオフセットとなります。一見すると、この可変長オフセットの計算は極めて重い処理に見えます。しかし、現代のプロセッサが備える高度なビット操作命令は、この計算を一瞬で消し去るマジックを秘めています。
6.2 Ternary reconstruction
6.2.1 0の復元
重みを復元する論理式は、極めてエレガントに記述されます。存在フラグを p ∈ {0, 1}、復元された符号フラグを n ∈ {0, 1}(負のとき1)と定義しましょう。
もし要素が不在(p = 0)であれば、符号フラグ n の値が何であろうとも、出力は厳密に 0 でなければなりません。
6.2.2 -1/+1の復元
もし要素が存在(p = 1)していれば、n = 0 のとき +1、n = 1 のとき -1 を出力しなければなりません。これを単一の代数式として統合したのが、原論文の式(2)です。
wi = pi - 2 × ni
検証してみましょう。
- 不在の場合(p=0, n=0): w = 0 - 0 = 0
- 正の存在の場合(p=1, n=0): w = 1 - 0 = +1
- 負の存在の場合(p=1, n=1): w = 1 - 2 = -1
条件分岐(if-else)など一行も必要ありません。ブール代数と単純な整数演算の組み合わせだけで、三元の宇宙が完全に復元されるのです。
6.2.3 Index mapping
この復元処理をSIMDレジスタの全レーンに対して並列に適用するためには、ビットマップの各ビットと、符号ストリームの各ビットを、レジスタ内の対応するバイト位置へと一挙に射影(Mapping)するアルゴリズムが要求されます。
6.3 なぜ普通の2-bit packingでは足りないのか
6.3.1 2-bit ternary
ここで改めて、業界で広く使われている標準的な2ビットパッキング方式と比較してみましょう。2ビットパッキングでは、例えば 00 = 0, 01 = +1, 11 = -1 といった固定のマッピングを用い、1バイトに4つの三元重みを均等に詰め込みます。
この方式の最大の利点は、デコードの単純さにあります。テーブル参照や単純なシフトとマスクだけで各重みを容易に抽出できます。しかし、ゼロ密度がどれほど高かろうと、メモリフットプリントは冷酷に 2.000 bits/weight に固定されます。BITCOSが 1.500 bits/weight を達成している横で、2ビットパッキングは33%も余計なトラフィックをメモリバスに垂れ流し続けることになります。
6.3.2 5-trit packing
一方の5-trit方式(llama.cppのTQ1_0など)は、1.625 bits/weight まで迫りますが、1バイトの中に含まれる数値を復元するために、乗算や除算、あるいは巨大な検索テーブルを必要とします。除算命令はプロセッサパイプラインにおいて最も重い命令の一つであり、これを内積ループの中で回すことは性能に対する自殺行為です。
6.3.3 BITCOS
BITCOSは、これら両者の「いいとこ取り」を狙った設計です。
- 空間効率においては、ゼロ密度を活用して5-tritを凌駕する 2 - z bits/weight を叩き出す。
- 計算効率においては、除算を一切排除し、ビット並列命令によって2-bitパッキングに迫る軽快さで復元を行う。
6.4 レイアウトはアルゴリズムである
6.4.1 Memory layout
BITCOSが我々に突きつける最も重要な教訓は、「データレイアウトとは、それ自体が計算アルゴリズムの具現化である」という認識です。データをどのように並べるかという決定は、受動的な保管庫の設計ではありません。それは、後続するプロセッサのデコード回路が、どのような順序でトランジスタをスイッチさせるかを直接プログラミングしていることに他なりません。
6.4.2 Cache behavior
存在ビットマップを独立したプレーンとして分離したことにより、プロセッサは「どの重みが非ゼロであるか」というメタデータのみを、極めて小さなメモリ転送量でL1キャッシュへと事前ロード(Prefetch)することが可能になりました。
6.4.3 Vectorization
ビットマップの1ワード(32ビット)が、そのままAVX-512の32レーンマスクレジスタ k1 へと直結します。データ構造の境界が、命令セットのレジスタ幅と分子レベルで一致しているのです。
6.4.4 Coalescing
GPU環境においても、16個の連続する出力チャネルを担当するワープ内の16スレッドが、隣接するビットマップワードを一括してコアレスド(合体)ロードできる配置が採用されています。美しきレイアウトは、美しき命令シーケンスを呼び寄せる磁石となるのです。
コラム:白紙のメモリマップ――構造体に魂を売った日
オブジェクト指向プログラミングに慣れ親しんだ若いエンジニアにコードレビューをすると、よく struct Weight { int8_t value; bool is_zero; }; のような構造体の配列(Array of Structures: AoS)を見かける。「可読性が高いですから!」と彼らは屈託なく笑う。筆者は泣きながらそのコードを全消しし、ビットマップと符号列を完全に引き裂いた構造(Structure of Arrays: SoA)へと書き直させる。メモリの上でデータを美しくカプセル化してはならない。データは、シリコンのベルトコンベアが最も掴みやすいように、解体され、平たく延ばされ、バラバラのビットストリームとして配置されなければならないのだ。プログラムの美学とは、抽象化の彼方ではなく、バス幅との完全な調和の中に宿る。
第7章 PDEPの錬金術――一命令が論文を変えるとき
7.1 Bit Manipulation
7.1.1 PDEP
BITCOSの理論を、x86プロセッサ上で実際に光速の実行コードへと昇華させた決定的な特効薬。それこそが、IntelがHaswellマイクロアーキテクチャ(2013年)で導入したBMI2命令セットに含まれる、ある風変わりな機械語命令――PDEP (Parallel Bits Deposit: 並列ビットデポジット) でした。
PDEP 命令の動作原理を図4に基づいて紐解きましょう。この命令は、2つの汎用レジスタを入力として取ります。ソースレジスタ a と、マスクレジスタ mask です。
dst = _pdep_u32(a, mask);
その動作は、擬似コードで書けば驚くほどシンプルです。ソース a の下位ビットから順番にビットを取り出し、マスク mask の中で「ビット1が立っている位置」に順番に落とし込んでいきます(Deposit)。そして、マスクがゼロである位置には、容赦なくゼロを書き込みます。
mask: 1 0 1 1 0 0 1 0 (存在ビットマップ: 4つの非ゼロ) a: 1 0 1 1 - - - - (コンパクト符号列: 4ビット) ------------------------------------------------------------ dst: 1 0 1 1 0 0 1 0 (元のレーン順に散布された符号列!)
刮目してください。この命令が実行している変換は、まさに「コンパクトに詰め込まれた符号ストリームを、存在ビットマップの指示に従って、元のテンソルの絶対座標へと一挙に散布(Scatter)する処理」そのものではありませんか!
7.1.2 Bit gather/scatter
通常、このような可変長のビット散布処理をソフトウェアのループで実装しようとすれば、ビットのテスト、条件分岐、可変シフト、OR結合といった命令を要素ごとに繰り返す必要があり、32要素を展開するだけでも数十から百サイクル以上の命令ストームを引き起こします。
ところが、Intelプロセッサの深部には、この処理を専用のクロスバー回路を用いてわずか1サイクル(近年のコアでは3サイクル程度)の単一命令で完了させるハードウェアが焼き付けられていたのです。
7.1.3 Prefix population count
符号ストリームから何ビットを消費したかは、ビットマップ自身のビット立脚数を数える単一の POPCNT(Population Count)命令で即座に判明します。ポインタの更新も一瞬です。かつてこれほどまでに、ある特異なデータ構造のために誂えられたかのような命令が存在したでしょうか。
7.2 命令レベルのボトルネック
7.2.1 Latency
Georganasらが書き上げたAVX-512向けアンパックシーケンス(図3)は、まさにマイクロアーキテクチャの極限芸術です。1回の内積ループ(32行分の重み処理)は、合計わずか17命令で構成されています。
重みの解凍そのものに直接関与する命令は、奇跡的なことにわずか3命令に凝縮されています。
pdep: 符号ビットをレーン順へ散布する。vmovdqu16 zmm5{k1}{z}: ビットマップマスク k1 の下で、ベーススケール +s をゼロマスク展開する。vmovdqu16 zmm5{k1}: 散布された符号マスクの下で、反転スケール -s を上書きマージする。
残りの14命令は、アドレス計算、アライメントのための可変シフト shrx、次のイテレーションのためのソフトウェアプリフェッチ、そして活性化ベクトルとの融合積和演算 vfmadd231ph です。
7.2.2 Throughput
このシーケンスにおいて、命令間の依存関係(Data Dependency Chain)は極限まで短縮されています。pdep は汎用整数パイプライン(Port 1等)で実行され、ベクトルの展開と積和演算はベクトルパイプライン(Port 0/5)で実行されるため、プロセッサのスーパースカラー実行エンジンは、整数処理と浮動小数点処理を完全に並行してインターリーブ(交互配置)させることができます。
7.2.3 Port contention
特定の実行ポートへの偏り(Port Contention)を徹底的に排除したこの配置により、ループの1イテレーションは、L1キャッシュ常駐時においてわずか 4.60サイクル(Emerald RapidsのPコア) という驚異的なクロック数で駆け抜けます。
7.2.4 Front-end pressure
ループ全体のコードサイズが極めてコンパクトであるため、プロセッサの命令デコードフロントエンドに負荷をかけることもありません。コード全体がプロセッサ内部のμOPキャッシュ(Decoded Stream Buffer: DSB)に完全に収まり、デコーダの電力消費すらカットしながらループが回転し続けるのです。
7.3 PDEPを発明ではなく「実装」に戻す
7.3.1 Representation algebra
しかし、ここで立ち止まって深呼吸をしなければなりません。敵対的査読者の批判を思い出してください。「それはIntelのPDEPというローカルな特効薬に依存したハックではないのか?」
この批判を乗り越えるために、我々は PDEP という特定の命令名を一度忘れ去り、この処理が数学的に何を意味しているのかを抽象化しなければなりません。我々が必要としているのは、特定のシリコン回路ではなく、以下の4つのプリミティブからなる「表現代数(Representation Algebra)」です。
- Extract: ビットマップから局所的な存在パターンを抽出する。
- Rank: 各要素の局所的な順位をプレフィックスサムで決定する。
- Select: 順位に基づいて符号ストリームから対応ビットを選択する。
- Permute: 選択された符号と存在フラグを、出力ベクトルの幾何空間へと再配置する。
7.3.2 Primitive operations
PDEP とは、この4段階の抽象代数を、x86がたまたま単一命令で融合実装していた具現化の一つに過ぎません。ならば、他のプロセッサ上であっても、これら4つのプリミティブをそのプロセッサが得意とする別の命令シーケンスで紡ぎ出せば、全く等価な表現代数が成立するはずです。
7.3.3 ISA mapping
事実、AVX-512のようなマスクレジスタを持たないクライアントCPU(AVX2環境)に対して、著者らは図5に示す全く別のシーケンスを構築しました。そこでは、マスクレジスタの代わりにバイトマスクを展開し、数学的関係式 wi = pi - 2ni を用いてAVX-VNNIの整数内積命令へと接続しています。特定の命令への依存から、表現代数のポータブルな写像へ。この視点の飛躍こそが、局所ハックを普遍理論へと変える架け橋となります。
7.4 Intel Xe2 SLMの教訓
7.4.1 Local memory
この「代数の別実装」の最も鮮やかな実例が、Intel Xe2 GPU向けに開発されたアンパックカーネルです(図6)。
GPUには、x86のような PDEP 命令は影も形もありません。そこで著者らは、GPU内部の超高速なオンチップ共有メモリ――共有ローカルメモリ(SLM: Shared Local Memory)に着目しました。
7.4.2 Shared execution
彼らは、4行分の存在ビット(4ビットのニブル)と、符号ストリームの次の4ビットを連結した「8ビットのキー」を考案しました。8ビットが取り得る状態数は 28 = 256 通りに過ぎません。この256通りのキーに対して、あらかじめ4行分の復元された三元FP16定数(4 × 2バイト = 8バイト)をテーブルとして計算しておきます。テーブル全体のサイズは、わずか 256 × 8バイト = 2 KB です。
2KBという極小のルックアップテーブル(LUT)は、各コンピュートユニットのSLMの中に余裕で収まります。カーネル起動時、ワークグループ内の各スレッドが協調してこの2KBのテーブルを一瞬でSLMへ展開します。
あとは、4行ごとにビットマップと符号列から8ビットのオフセットを合成し、単一の load.slm.d32x2 命令を発行するだけです。1回のテーブル参照で、4行 × 16列 = 64個のFP16三元重みが、完全にVNNI2の積和演算順序に整列した状態でレジスタへと雪崩れ込み、直後のDPAS(Dot Product Accumulate Systolic: XMXシストリックアレイ)命令へと直結されます。
7.4.3 「速い実装」と「一般理論」の分離
CPUではハードウェアビット散布命令(PDEP)を使い、GPUでは共有メモリ内の事前計算テーブル(SLM LUT)を使う。表層の機械語命令は似ても似つきません。しかし、背後で機能している数学的本質――「存在ビットマップの導きによって、沈黙の空間をスキップしながら符号を復元する」という表現代数は、完全に同一です。速い実装の個別性に惑わされず、一般理論の普遍性を抽出すること。これこそが、次章以降で展開されるPATE理論の揺るぎない土台となるのです。
コラム:消えた命令を探して――シリコンアーキテクトの悪戯
命令セットマニュアルの片隅に、用途不明の奇妙な命令がひっそりと記載されていることがある。BMI2の PDEP と PEXT(パラレルビット抽出)を最初に見つけたとき、大抵のプログラマは「チェスのビットボード(盤面表現)を高速化するためだけにIntelのエンジニアが悪ノリで作った珍命令」だと笑っていた。実際、長年その用途以外で日の目を見ることは滅多になかった。まさか10年以上が経過した後、世界を揺るがす巨大言語モデルの推論を救う中核兵器としてこの命令が脚光を浴びることになろうとは、設計したアーキテクト自身すら予期していなかったに違いない。シリコンの中に埋め込まれた種子は、全く異なる時代のアルゴリズムと出会ったとき、突然狂い咲くように花を咲かせるのだ。
第8章 ゼロを運ぶ――圧縮データのMemory Hierarchy
8.1 DRAMからCacheへ
8.1.1 Payload削減
命令レベルの錬金術によってアンパックの目処が立ったとき、我々の視線は再びプロセッサの広大なメモリ階層全体へと向けられます。BITCOSがもたらす最大の恩恵は、何と言ってもDRAMから吸い上げる物理的なペイロード(転送データ量)の劇的な削減です。
ゼロ密度 z = 0.45 のモデルを例にとれば、実効ビット幅は 1.55 bits/weight となり、従来の16ビット浮動小数点表現に対して 10.3分の1、標準的な2ビットパッキングに対しても 約1.29分の1(22.5%削減) のトラフィック削減を達成します。メモリバスの上を飛び交うトランザクションが2割以上減少するということは、同一のメモリ帯域幅の下で、原理的に1.29倍の反復速度で重みを流し込めることを意味します。
8.1.2 Cache residency
さらに見落としてはならないのが、プロセッサ内部のキャッシュ階層に対する副次的効果です。近年のサーバー向けプロセッサは数十メガバイトから数百メガバイトの巨大なラストレベルキャッシュ(LLC: Last Level Cache)を搭載していますが、パラメータが数十億に及ぶLLMの前には、LLCすら瞬時に溢れかえります。
しかし、データサイズが2割削られることによって、従来はLLCに収まりきらずにDRAMへ溢れ出していた重要な重みタイルやアテンションのKVキャッシュの一部が、キャッシュの境界線の内側に踏みとどまる(Cache Residencyが向上する)という非線形な恩恵がもたらされます。
8.1.3 Cache miss
DRAMへのアクセスレイテンシ(約60〜80ナノ秒)と、L3キャッシュへのアクセスレイテンシ(約15〜20ナノ秒)の間には、4倍以上の速度差が存在します。ペイロードの線形な削減が、キャッシュミスの劇的な低減をトリガーし、実測の推論レイテンシを理論予測以上に引き下げる「キャッシュのボーナスステージ」が出現するのです。
8.2 圧縮されたまま読む
8.2.1 Decode-on-load
このメモリ階層の恩恵を最大化するために、推論エンジンはデータの解凍を「どこで」行うべきでしょうか。古典的なシステム設計では、メモリから読み出した圧縮データを一度別のバッファ領域へ解凍して展開し(Decode-on-load)、その展開済み配列に対して標準的なGEMM関数を呼び出すという二段階のアプローチが採られがちでした。
しかし、LLMのデコード推論においてこのアプローチは致命的な大失敗となります。解凍されたデータを再びキャッシュやメモリへ書き戻すトラフィックが発生し、貴重なメモリ帯域を自ら食いつぶしてしまうからです。
8.2.2 Decode-on-use
正解は、「Decode-on-use(使用時解凍)」、すなわち演算器が重みを消費するまさにその瞬間まで、データは圧縮された姿のままメモリ階層を駆け上がり、レジスタの直前、あるいはレジスタ内部においてのみ展開されるという究極のパイプライン融合です。
BITCOSのAVX-512シーケンスはこの哲学を徹底しています。圧縮されたビットマップと符号列は、DRAMからL3、L2、L1キャッシュを経て、そのままの姿で汎用レジスタへとロードされます。そして、FMA命令が発行される直前のわずか数クロックの間に pdep とマスク操作によって zmm レジスタ内でFP16へと復元され、次の瞬間には活性化ベクトルと掛け合わされてアキュムレータへと吸収されます。解凍された生データは、チップ上のどの配線にも、どのキャッシュラインにも、1ナノ秒たりとも痕跡を残しません。
8.2.3 Expand-once
ただし、このDecode-on-useの鉄則も、再利用係数 R が極めて大きいPrefill相においては、再びその妥当性を疑われることになります。同一の重みを数千回使い回すのであれば、最初の1回だけローカルSRAM(GPUのShared Memoryなど)に解凍・展開して保持し(Expand-once)、以降はその展開済みデータを全員でしゃぶり尽くす方が、解凍命令のトータルコストを削減できるからです。状況に応じたこの戦略の転換こそが、次章の主題となります。
8.3 圧縮と局所性のトレードオフ
8.3.1 Small payload
ペイロードを小さくすることの代償として、システムは常に局所性と複雑性のトレードオフという名の利息を支払わされています。
8.3.2 More instructions
ペイロードを削れば削るほど、ビットを解きほぐすための命令数は増加します。メモリ帯域が狭い環境では、この追加命令はメモリ待ち時間の背後に完全に隠蔽されますが、メモリ帯域が潤沢な環境では、余計な命令の山として露出し、コアの実行効率を蝕みます。
8.3.3 Fewer bytes, more work
「バイト数は減ったが、仕事量は増えた(Fewer bytes, more work)」。この冷厳な事実を認識することなしに、高性能コンピューティングの海を航海することはできません。
8.4 Energyという第四の軸
8.4.1 pJ/bit
そして、速度(Latency/Throughput)の議論の背後に控えているのが、現代の持続可能コンピューティングにおける究極の審判――エネルギー消費(Energy Consumption)です。
マーク・ホロウィッツ(Mark Horowitz)の記念碑的なエネルギー分析が示す通り、シリコンチップ上において32ビット整数の加算を実行するのに消費されるエネルギーはわずか 0.1 pJ(ピコジュール) 程度です。浮動小数点積和演算(FP32 FMA)であっても数pJに過ぎません。
それに対して、DRAMからチップ外の配線を通じてデータを1ビット読み出すのに消費されるエネルギーは、実に 数十から数百 pJ/bit に跳ね上がります。プロセッサが消費する電力の大部分は、演算そのものではなく、ビットを空間的に移動させること(Data Movement)のために費やされているのです。
8.4.2 pJ/MAC
BITCOSがもたらす20%以上のペイロード削減は、単に推論トークン速度を向上させるだけではありません。それは、トークンを1つ生成するごとにDRAMの充放電によって大気中へと熱として捨てられていた莫大なジュール数を、物理的に削減していることを意味します。
8.4.3 Edge inference
バッテリーで駆動するスマートフォン、熱冷却が厳しく制限されたドローンや車載カメラにおいて、この「1トークンあたりのエネルギー(Joules per Token)」の削減は、デバイスの動作継続時間と熱暴走マージンを決定づける死活問題です。ゼロを巧みに運び、無駄な転送を根絶すること。それは、速度の追求であると同時に、物理的な限界に縛られた現実世界に対する、計算機工学の最も誠実な応答なのです。
コラム:熱風のサーバルーム――電気代という名の現実
真夏のデータセンターのサーバルームに足を踏み入れたことがあるだろうか。轟音を立てて回転する冷却ファンと、ラックの背面から吐き出される目眩のするような熱風の壁。吸気口と排気口の温度差を見つめながら、筆者はいつも計算の「質量」を感じる。クラウドの彼方にある知能とは、クリーンな数学の結晶などではない。莫大な火力発電と原子力発電の電力を貪り食い、熱へと変換し続ける巨大な抵抗器なのだ。三元モデルのビットを削り、DRAMのバスの点滅を1回でも減らすこと。それは、ディスプレイの前のプログラマが、地球の物理的な限界に対して行使できる数少ない抵抗の一つなのだと信じている。
第三部 表現の相転移――PATE理論とCritical-Path-Hidden Unpacking
第9章 PATE――重み表現は「モデル属性」ではない
9.1 Representation as Execution Policy
9.1.1 Static formatという前提
現代の機械学習工学における最大の盲点は、「モデルの重み表現(Weight Representation)とは、訓練完了時に決定され、ディスク上に直列化(Serialize)された不変の静的アーティファクトである」という頑固な思い込みにあります。Hugging FaceのリポジトリからSafetensorsファイルをダウンロードし、それをVRAMやメインメモリにロードした瞬間から、テンソルは単一の固定されたビット幅――FP16、INT4、あるいは三元モデルにおける1.58ビットや2ビット――として凍結されます。従来のサービングフレームワークは、この固定されたビット列をいかに速く演算器へと送り届けるかという「受動的な配管作業」に終始してきました。
9.1.2 Runtime representation
しかし、計算機システムの冷厳な現実に照らし合わせるならば、この前提は根本から覆されなければなりません。プロセッサにとって、データフォーマットとは単なる保存形式ではなく、命令パイプラインの振る舞い、キャッシュコヒーレンシの維持コスト、そしてメモリバスの占有率を直接支配する「動的なシステムリソースの配分方針」そのものです。
我々はここで、新しいパラダイムを宣言します。重み表現とは、モデル固有の属性(Static Model Attribute)ではなく、実行基盤(ランタイム)が動的な状況変化に応じて柔軟に選択・変更すべき「実行時ポリシー(Runtime Execution Policy)」である、と。
9.1.3 Hardware-aware inference
ハードウェアの物理構造、利用可能なメモリ帯域幅、コア数、そして推論のバッチサイズ。これら刻々と変動するコンテキストに応じて、重みテンソルはある瞬間にはギチギチに詰め込まれたストリームとなり、次の瞬間には演算器が最も喜ぶ非圧縮のタイルへと姿を変えるべきです。この哲学を定式化した体系こそが、PATE (Phase-Adaptive Ternary Execution: 相適応型三元実行) 理論です。
9.2 四つの状態変数
PATE理論において、システムが取り得る最適な重み表現 F は、真空の中に孤立して存在するのではなく、厳密に定義された4つの状態変数(State Variables)によって構成される状態空間上の最適解として決定されます。
9.2.1 Zero density (z)
第一の変数は、重みテンソル固有の統計的疎性を示すゼロ密度 z ∈ [0, 1] です。第5章で解剖した通り、この値が 0.375 を上回るか否かによって、BITCOSのような疎性適応型レイアウトが固定長パッキングに対して理論的優位を獲得できるかどうかが決まります。
9.2.2 Arithmetic intensity (AI)
第二の変数は、ワークロードの処理特性を規定する算術強度 AI [FLOPs/Byte] です。1バイトの重みをメモリから読み出す間に、システムが何回の計算を行えるか。後述するように、この値はDecode相とPrefill相の間で数桁にわたって激しく乱高下します。
9.2.3 Machine balance (B/P)
第三の変数は、ハードウェアの物理的骨格を規定するマシンバランス ρ = B / P [Bytes/FLOP] です。ここで B は持続可能メモリ帯域幅 [Bytes/s]、P はプロセッサの実効演算性能 [FLOPs/s] です。この比率は、プロセッサが1回の演算を行う間に、メモリから何バイトのデータを供給できるかという「配管の太さ」を示しています。
9.2.4 Reuse (R)
第四の変数は、メモリから一度フェッチされた重みデータが、破棄されるまでに何回繰り返し演算に供されるかを示す重み再利用度 R [Uses/Fetch] です。バッチサイズが1であれば R ≈ 1 ですが、Prefill相や大バッチ推論においては R は数十から数千へと膨れ上がります。
9.3 PATEの中心式
9.3.1 最適表現選択の定式化
これら4つの変数を統合し、PATE理論は推論時における最適表現 F* の決定問題を、次の最小化問題として定式化します。
F* = argminF ∈ F T( F ; z, AI, ρ, R )
ここで候補集合 F は、高密度な5-tritパッキング、固定長の2-bitパッキング、疎性適応型のBITCOS、そして事前にローカルSRAMへ解凍された展開三元タイル(Expanded Ternary Tile)など、互いに異なるトレードオフを持つ表現ポートフォリオ(Representation Portfolio)です。
9.3.2 Memory term
表現 F を採用した際、重み行列をメモリ階層間で移動させるのに要する時間 Tmemory は、実効ビット幅 bF(z) と再利用度 R の関数として定義されます。
Tmemory(F) = ( bF(z) × Nw ) / ( B × freuse(R) )
9.3.3 Compute term
一方、その表現を用いて演算器を駆動する時間 Tcompute は、有用な積和演算量 W と、その表現に最適化された計算カーネルの実効スループット PF によって決まります。
Tcompute(F) = W / PF
9.3.4 Unpacking term
そして決定的なのが、圧縮された表現 F を演算可能なレジスタ形式へと復元するためのアンパック時間 Tunpack(F) です。従来のシステム工学は、この項を単純に総実行時間に加算していました。しかしPATE理論は、次章で詳述するパイプラインスラック(Slack)への重複隠蔽を導入することにより、このアンパック項を条件付きで完全にゼロ化する数理的枠組みを提示します。
9.4 「最適表現」は一つではない
9.4.1 Decode representation
この中心式から導き出される直接の結論は、推論システムにおける「銀の弾丸(万能のフォーマット)」の完全な否定です。自己回帰デコードのように R ≈ 1 であり、AI < ρ となる極端なメモリ帯域律速下では、Tmemory を最小化するBITCOSのような高圧縮表現が圧倒的な最適解となります。
9.4.2 Prefill representation
対照的に、Prefill相のように R ≫ 1 であり、AI > ρ となる計算律速下では、Tunpack を一度の展開で償却(Amortize)し、PF を最大化できる非圧縮タイル展開表現が玉座を奪取します。
9.4.3 Batched representation
同時リクエストが増大するバッチ推論では、バッチサイズ N の増加に伴って最適な表現形式が動的にシフトします。
9.4.4 Edge representation
そしてリソースが極度に制約されたエッジデバイスでは、メモリフットプリントと消費エネルギーを最小化する極限のバランス点が選択されます。重み表現を固定化する時代は終わりました。PATEは、単一の静的チェックポイントから、状況に応じて無限の身体を編み出すための理論的コンパスなのです。
コラム:カメレオンの皮――変幻自在なデータの身体
動物行動学のドキュメンタリーで、周囲の岩肌に合わせて皮膚のテクスチャと色彩を瞬時に変化させるタコを見たときの衝撃を今でも覚えている。あれこそが究極のアーキテクチャだ。捕食者(メモリウォール)が迫れば極限まで小さく身を縮め、獲物(大量のバッチトークン)が飛び込んできたら巨大な触手を広げて一挙に飲み込む。なぜ我々は、半導体の上で動くデータに対して「お前はずっと直方体のブロックのままでいろ」と強要してきたのだろうか。データに筋肉と皮膚を与え、シリコンの地形に合わせて変幻自在に波打たせること。PATEの数式をホワイトボードに走らせながら、我々はデータ構造に生命を吹き込んでいるのだという奇妙な高揚感に包まれていた。
第10章 相転移――DecodeからPrefillへ、同じ重みが別の姿になる
10.1 Decode Regime
10.1.1 Batch=1
LLM推論の現場において、自己回帰デコード(Decode Phase)は「砂漠を往く隊商」に喩えられます。1つのトークンという極小の荷物を運ぶために、数ギガバイトから数十ギガバイトに及ぶ重み行列の全量を、メモリの彼方から毎回引きずり出さなければなりません。
10.1.2 Weight streaming
このレジームにおいて、重みデータはメモリからプロセッサへと高速に通り過ぎていくだけの「使い捨てのストリーム(Streaming Data)」です。重みがキャッシュに留まる暇はなく、バスの上を流れるビット幅の細さそのものが、トークン生成速度の絶対的な上限を画定します。
10.1.3 Bandwidth-bound execution
ここでは、演算器がいくら高性能であろうと関係ありません。メモリバスが1秒間に運べるバイト数がすべてを支配します。したがって、Decodeレジームにおける唯一至高の戦略は、「1ビットでも細くしてメモリバスを通過させること」、すなわち圧縮ストリーミング実行(Compressed-Stream Execution)です。
10.2 Prefill Regime
10.2.1 Token parallelism
ところが、ユーザーが長大なコンテキストを入力し、システムがプロンプトの理解を開始した瞬間、推論の物理世界は劇的な相転移を起こします。これがPrefillフェーズです。数千のトークンが同時に行列として入力されるため、計算の幾何学はベクトル積(GEMV)から密な行列積(GEMM)へと跳躍します。
10.2.2 Reuse
重み行列の各ブロックは、一度オンチップのキャッシュやSRAMに読み込まれれば、数千個の入力トークンベクトルに対して繰り返し積和演算を実行するために再利用されます。再利用係数 R は容易に 1000 や 2048 を超えます。
10.2.3 Compute-bound execution
この状況下では、メモリバスの帯域幅はもはや律速段階ではありません。プロセッサ内部のシストリックアレイやSIMD積和演算ユニットが、1秒間に何千億回、何兆回の足し算を実行できるかという「計算律速(Compute-bound)」へと支配権が完全に移行します。
10.3 Representation Phase Transition
10.3.1 Roofline crossing
この二つのレジームの境界線上で、極めてスリリングな現象が発生します。ルーフラインモデルのグラフ上において、ワークロードの算術強度 AI が、プロセッサのマシンバランス ρ = B / P と交差する瞬間です。
AI(R) = ( Useful FLOPs ) / ( Bytes Fetched ) ∝ R
再利用度 R の増大に伴って、動作点はルーフラインの斜面(メモリ律速域)を駆け上がり、平坦な天井(計算律速域)へと突入します。
10.3.2 Reuse threshold
この交差点において、最適な重み表現の選択が不連続に切り替わる臨界点――表現相転移(Representation Phase Transition)の閾値 R* が存在します。
R < R* : Compressed-Stream Mode (BITCOS等の極小ビット幅が勝利)
R > R* : Tile-Expanded Mode (INT8事前展開タイルが勝利)
10.3.3 Representation crossover
もしPrefillフェーズにおいて、Decode用のBITCOS圧縮ストリームをそのまま使い続けようとすれば何が起きるでしょうか。重みが使われるたびに、毎回 PDEP やSLMルックアップテーブルを叩いて解凍処理が繰り返され、演算器の手前で無駄なアンパック命令が山積します。
逆に、Tile-Expanded Modeを採用すれば、最初のフェッチ時に一度だけ重みタイルをレジスタや共有メモリ上で非圧縮の三元整数(INT8形式の -1, 0, +1)へと展開し、以降の数千回の積和演算はハードウェアが最も得意とするネイティブなシストリック積和命令(Intel AMXやNVIDIA Tensor Core MMA)にノーガードで叩き込ませることができます。解凍にかかった初期コストは、巨大な R によって完全に希釈され、極限までゼロに近づきます。
10.4 一つのモデル、複数の物理的身体
10.4.1 Compressed stream
静的な思考に囚われたエンジニアは問うでしょう。「では、ディスクやメモリにはどちらの形式で保存しておくべきなのか?」と。
10.4.2 Expanded tile
PATEの回答は明快です。「主記憶(DRAM)上には最もコンパクトな圧縮ストリーム形式(BITCOS)でただ一つだけ保持し、メモリ階層を駆け上がるプロセスの途上で、ランタイムが相転移を引き起こす」のです。
10.4.3 Hybrid representation
DRAM上にある単一の三元チェックポイントが、Decode時には圧縮された毛細血管のようなストリームとしてコアへと注ぎ込まれ、Prefill時にはL2キャッシュや共有メモリのプールの中で瞬時に巨大な非圧縮タイルへと受肉する。一つの知能モデルが、実行フェーズに応じて異なる複数の物理的身体を使い分ける。これこそが、相適応型実行が切り拓く新しい推論の地平です。
コラム:物質の三態――氷が水になり、蒸気となるように
高校の物理で、氷に熱を加えていくと温度上昇がいったん止まり、水へと相転移する「潜熱」のグラフを習ったはずだ。状態が変わるとき、そこには必ずエネルギーの不連続な出入りがある。LLMの推論を観察していて、あのグラフと全く同じ光景を見たときの戦慄は忘れられない。プロンプトのトークン長が64を超えた瞬間、プロファイラのパイプラインストール要因が「メモリデータ待ち」から「実行ポート競合」へと垂直落下するように切り替わったのだ。データはもはや固体(固たく詰め込まれた圧縮ビット)ではいられず、液体(キャッシュを循環する展開タイル)へと融解することを求めていた。自然界の物理現象は、いつだってシリコンの回路の中で密かに繰り返されている。
第11章 消えるアンパック――Critical-Path-Hidden Unpacking
11.1 「ゼロコスト」という危険な言葉
11.1.1 命令は消えない
計算機アーキテクチャの論文において、「オーバーヘッドはゼロである(Zero-cost)」という文句ほど、査読者を激怒させ、そして疑心暗鬼に陥れる言葉はありません。熱力学の第二法則が教える通り、この宇宙において仕事を行うためには必ずエネルギーと時間が必要です。解凍のためのビットシフト命令も、マスク生成命令も、物理的なトランジスタを充放電させ、実行パイプラインのステージを確実に通過しています。命令そのものが魔法のように消失することなど絶対にあり得ません。
11.1.2 latencyとcritical path
しかし、Out-of-Order(アウト・オブ・オーダー: 命令の追い越し実行)実行エンジンと、高度に並列化されたメモリサブシステムを備える現代のプロセッサにおいては、「命令が物理的に実行されていること」と「その命令のレイテンシが全体の完了時間を遅延させること」は全く同義ではありません。
11.1.3 Visible overheadという概念
システム全体の実行時間を決定づけている最長の一連の依存関係の連鎖、すなわちクリティカルパス(Critical Path)。解凍命令がこのクリティカルパスの外側に押し出され、他の不可避な待ち時間の影に完全に隠れ果てているならば、外部の観測者にとって、そのアンパック処理のレイテンシは「不可視(Zero Visible Overhead)」となります。
11.2 Slack
11.2.1 Memory slack
この不可視化を可能にする物理的メカニズムが、実行パイプラインにおける「スラック(Slack: 隙間時間/弛み)」の存在です。
いま、メモリ帯域律速なDecode相(バッチ1)において、ある重みブロックをDRAMからL1キャッシュ経由で引き抜くのに要する時間を TM、その重みを用いて積和演算を実行するのに要する時間を TC としましょう。メモリウォールに囚われたこの領域では、TM は TC よりも圧倒的に長大です。
TM ≫ TC
演算器は、次の重みデータがDRAMから届くのを待つ間、完全に手持ち無沙汰な空白時間――すなわちメモリスラック SM = TM - TC を抱えて立ち往生しています。
11.2.2 Compute slack
逆に、計算律速なPrefill相においては、積和演算器が猛烈に回転している傍らで、メモリ転送バス側には巨大な余裕――コンピュートスラック SC = TC - TM が生まれます。
11.2.3 Pipeline slack
これらのスラックは、現代のプロセッサ内部に構造的に遍在する「失われた時間」のポケットです。
11.3 Zero-visible-overhead condition
11.3.1 定理の定式化
PATE理論は、このスラックの幾何学に基づき、アンパック処理が真に不可視化されるための必要十分条件を、極めて厳密な不等式として打ち立てます。これが「臨界経路隠蔽条件(Zero-Visible-Overhead Condition)」です。
アンパック処理そのものを実行するのに要する正味の時間を TU と置くとき、以下の不等式が満たされるならば、アンパック処理は全体の実行時間を1サイクルたりとも延長しません。
TU ≤ | TM - TC |
11.3.2 Sufficient condition
この条件式の美しさは、Decode相とPrefill相の双方を単一の不等式で貫通している点にあります。
- Decode相(TM > TC): スラック SM = TM - TC の内部に、アンパック命令列を完全にインターリーブ(挟み込み)して潜り込ませることができれば、TU ≤ SM となり、アンパック時間は長大なDRAMレイテンシの陰に完全に隠蔽されます。
- Prefill相(TC > TM): 演算器が過去のタイルをしゃぶり尽くしている膨大な時間 SC = TC - TM の裏で、非同期に次の圧縮タイルをフェッチして解凍を完了させておけば(ダブルバッファリング)、TU ≤ SC となり、演算器から見たアンパック待ちは厳密にゼロとなります。
11.3.3 Necessary conditionとの違い
「アンパックにはコストがない」と強弁することは科学的詐術です。しかし、「アンパック処理は、物理的なスラック不等式 TU ≤ |TM - TC| が成立する領域において、クリティカルパスから完全に消去されている」と主張することは、測定可能で反証可能な正統的計算機科学の命題です。この概念的転換によって、敵対的査読者の最大の武器であった「解凍オーバーヘッド批判」は、理論の前提条件へと鮮やかに解体されます。
11.4 Fused execution
11.4.1 Load
この隠蔽条件を具現化するための究極のソフトウェア実装が、完全融合実行カーネル(Fully Fused Execution Kernel)です。
11.4.2 Unpack
別個の解凍ループを回すのではなく、単一のタイトな内積ループの内部において、ロード命令、アンパック命令、そして演算命令が、アセンブリレベルで完璧な織物のようにインターリーブされます。
11.4.3 MMA/MAC
レジスタへのロードが発行された直後、メモリコントローラからのデータ返還を待つパイプラインの空きスロットに、直前のイテレーションで届いたデータの pdep やマスク操作が差し挟まれます。そして解凍が完了した瞬間、間髪を入れずにベクトルFMA命令やDPAS命令が同一レジスタを消費します。
11.4.4 Store
一時的な中間データは、プロセッサのレジスタファイルの外側へ一滴も漏れ出すことはありません。ロードの遅延、アンパックの計算、そして積和の集約。三者が互いの影に隠れ合いながら単一の呼吸で回転するこの融合カーネルこそが、スラックを極限まで食らい尽くす物理の錬金術なのです。
コラム:忍者の歩法――踏みしめる足音を消す技術
昔の忍術書を読むと、敵に忍び寄るときは「相手が息を吐く瞬間」や「風が木々を揺らす瞬間」に合わせて足を前に出せと書いてある。周囲の環境が立てる大きなノイズの影に、自分の小さな足音を完全に重ねて消し去るのだ。Critical-Path-Hidden Unpackingのコードを書いているとき、自分はシリコンの忍者なのだと思う。DRAMという鈍重な巨人が「ドシン」と音を立ててデータを運んでくるその巨大な足音の隙間に、ササッと数命令の pdep とシフトを滑り込ませる。巨人が顔を上げたときには、もうデータは何食わぬ顔でFP16に展開されてFMAの前に整列している。プロファイラを見ても、忍者の姿はどこにも映らない。ただ、結果としての速度だけがそこにある。
第12章 ARMはPDEPを必要としない――ISAから表現を解放する
12.1 Representation Algebra
12.1.1 Presence
第3章において提示された敵対的査読者の第三の異議――「PDEPというIntel特有の命令に依存した局所解ではないか」という批判に対し、我々は今こそ決定的な回答を与えなければなりません。
第7章で予告した通り、BITCOSの本質は特定の機械語命令のニーモニックにあるのではなく、存在ビットマップから三元テンソルを復元する「表現代数(Representation Algebra)」という抽象的数学構造にあります。
12.1.2 Rank
この代数は、いかなるプロセッサであれ、次の4つの不可欠なステップによって構成されます。
- Presence Extraction: ビットマップから局所的な非ゼロ存在フラグを取り出す。
- Rank Computation: プレフィックスサムによって、非ゼロ要素のコンパクトストリーム内での絶対順位を同定する。
- Sign Extraction: 順位に基づいて符号ストリームから該当ビットを読み出す。
- Vector Selection: 存在と符号を合成し、演算器のレーン幾何学へと射影する。
12.1.3 Sign
この4段階の数学的変換が保証される限り、ハードウェアが提供する命令が何であれ、アルゴリズムは等価に成立します。
12.1.4 Select
プロセッサの個性とは、この4段階の代数を、どの回路ユニットを用いて、いかに少ないサイクル数で駆け抜けるかという「表現の具現化手法」の違いに過ぎません。
12.2 SVE2
12.2.1 Predicate
では、世界中のモバイルデバイスや、AWS Graviton、富岳のようなスーパーコンピュータを支配するARMアーキテクチャにおいて、この代数はどのように具現化されるのでしょうか。最新のARMv9世代におけるベクトル拡張であるSVE2 (Scalable Vector Extension 2) を見てみましょう。
12.2.2 Vector bit manipulation
SVE2には、オプション機能拡張として FEAT_SVE_BitPerm が定義されています。そしてそこには、驚くべきことに BDEP (Bit Deposit) という命令が存在します。この命令は、SVEの可変長ベクトルレジスタの各64ビット要素内において、x86のPDEPと完全に同一のビット散布処理をベクトル並列で実行します。
bdep z2.d, z0.d, z1.d
SVE2が利用可能な環境であれば、x86のAVX-512シーケンスとほぼ1対1に対応するエレガントなベクトルコードを直接書き下すことが可能です。
12.2.3 Gather/compact
さらにSVE2は、強力な述語レジスタ(Predicate Registers)と、それに基づく COMPACT や SPLICE といった先鋭的なベクトル圧縮・展開命令を備えています。ビットマップを直接述語レジスタへと変換し、レーン単位のギャザー/スキャッターを条件分岐なしで流し込む表現力において、SVE2はx86を凌駕するポテンシャルを秘めています。
12.2.4 Scalable vector length
ハードウェアのベクトル長が128ビットであれ512ビットであれ、同一のバイナリが無修正でスケールして実行されるというSVE2の柔軟性は、PATEの適応型ポリシーにとって最も親和性の高いゆりかごとなります。
12.3 NEON
12.3.1 Lookup
しかし、世の中に無数に出回っている既存のスマートフォンや低消費電力SoCは、SVE2を持たない標準的な128ビット幅のNEONエンジンで駆動しています。NEONには BDEP のような気の利いたビット操作命令はありません。ここでBITCOSは座礁するのでしょうか?
否です。我々には、NEONが伝統的に極めて得意とする武器――ベクトルテーブル参照命令 tbl (Table Vector Lookup) があります。
12.3.2 Shift/mask
Intel Xe2 GPUがSLMで行った工夫を思い出してください。4ビットの存在ニブルと4ビットの符号ニブルを結合した8ビットのキーを用いれば、4要素分の復元定数は高々256エントリのテーブルに収まります。NEONの tbl 命令は、レジスタ内に配置されたバイトテーブルから、インデックスレジスタの値に基づいて16バイトのデータを瞬時に並列抽出することができます。
12.3.3 Fixed-width tile decode
固定幅の4ビットチャンクごとにシフトと論理積でキーを合成し、tbl で一挙に三元ベクトルを展開する。ハードウェアに専用の散布回路がないならば、小さなルックアップテーブルをレジスタ内に構築して迂回する。この「表現代数のNEONマッピング」によって、BITCOSはARMの世界においても完全に実用的なアンパックスループットを達成します。
12.4 「等価」とは何か
12.4.1 命令の等価性ではない
この事実は、ソフトウェア移植性(Portability)の概念を根底から再定義します。真の可搬性とは、異なるプロセッサ間で「同一の機械語命令が実行できること」ではありません。
12.4.2 操作意味論の等価性
それは、異なる命令セットアーキテクチャが提供する独自のプリミティブを組み合わせて、「同一の操作的意味論(Operational Semantics)を持つ表現代数を構築できること」です。
12.4.3 Critical-pathの等価性
そして、その代数がそれぞれのハードウェアのスラック方程式を満たし、クリティカルパスからアンパック遅延を等しく消去できていること。命令の類似性に惑わされるのをやめ、代数の等価性に立脚したとき、BITCOSとPATE理論は、Intelという特定のクレードル(揺り籠)を脱ぎ捨てて、世界中のあらゆるシリコンの上へと解き放たれるのです。
コラム:バベルの塔のプログラマ――方言を愛するということ
x86のアセンブラを愛するエンジニアと、ARMのアセンブラを愛するエンジニアを同じ部屋に閉じ込めると、大抵ろくなことにならない。「可変長命令の美しさがわからんのか」「ロード・ストアアーキテクチャの直交性こそが正義だ」と、不毛な宗教裁判が始まる。だが、BITCOSの代数をARM NEONの tbl 命令の上に美しくマッピングできた深夜、筆者は両方のアーキテクチャに対して深い感謝の念を抱いた。彼らは異なる方言を喋っているだけなのだ。表現しようとしている数学の詩が同じであるならば、単語の違いなど些細な問題に過ぎない。バベルの塔の崩壊を嘆くのではなく、多様な方言が織りなす響きの豊かさを楽しむ境地に達したとき、エンジニアはプラットフォームの奴隷から、真のシステムアーキテクトへと脱皮する。
第13章 Tensor Coreの前では、ゼロも行列になる
13.1 GPUはCPUと違う
13.1.1 Warp
思考の実験場を、再び超並列の巨獣――モダンGPUの世界へと移しましょう。GPUの実行モデルは、CPUのそれとは根本的に異なる物理法則に支配されています。GPUは、単一の強力なスレッドが複雑な分岐をこなす場所ではありません。32本のスレッドが完全に足並みを揃えて同一の命令を実行するワープ(Warp: NVIDIA)、あるいは16スレッドからなるサブグループ(Subgroup: Intel)が、何百、何千と集積した並列性のるつぼです。
13.1.2 Shared Memory
各コンピュートユニットの内部には、スレッド間でデータを共有するための極めて広帯域かつ低レイテンシなオンチップメモリ――Shared Memory(NVIDIA)/SLM(Intel)が備わっています。
13.1.3 Register File
そして各スレッドには、膨大な数の物理レジスタが与えられています。しかし、この豊かなリソースの背後には、スレッド間の「ダイバージェンス(分岐の不揃い)」や「共有メモリのバンク衝突(Bank Conflict)」が発生した瞬間に、性能が10分の1に急落するという過酷な地雷原が広がっています。
13.2 Compressed Global Memory
13.2.1 Load
GPU環境においてBITCOSをスケールさせるための鍵は、メモリコントローラの要求する合体アクセス(Coalesced Access)を維持しながら、圧縮データをいかにロードするかという「幾何学の整合」にあります。
グローバルメモリ(HBMやGDDR)上において、存在ビットマップは出力チャネル方向に連続して配置されます。これにより、ワープ内の連続するスレッド群は、単一の幅広トランザクションによって、隣接する列のビットマップワードを一括して引き抜くことができます。
13.2.2 Cooperative unpack
しかし、符号ストリームは各列の非ゼロ数に応じて不規則なオフセットを持ちます。ここで個々のスレッドが勝手気ままに符号メモリへアクセスすれば、非合体アクセスの嵐によってメモリパイプラインは崩壊します。
解決策は、ワープ協調型アンパック(Cooperative Warp Unpacking)です。ワープ内のスレッド群が、互いの符号オフセットをワープシャッフル命令(__shfl_sync)やプレフィックスサム命令(__popc と __ballot_sync の組み合わせ)を用いて共有し、符号ストリームの読み出しを整然としたブロック転送へとスケジューリングします。
13.2.3 Tile formation
引き出された圧縮データは、共有メモリ上に一時的な2次元タイルとして形成されます。
13.3 Unpack→MMA
13.3.1 Register tile
そして、最も肝要なステージが訪れます。展開された三元重みを、GPUの誇る最強の行列計算エンジン――NVIDIAのTensor Core、あるいはIntelのXMX (Xe Matrix eXtensions) へと受け渡すプロセスです。
13.3.2 Tensor Core
現行世代のTensor Core(Hopperのwgmma命令やBlackwellのtcgen05命令群)は、三元シンボルを直接入力として受け付けるようには作られていません。彼らが貪り食うのは、FP16、BF16、あるいはINT8やFP4といった、厳密に規格化された低精度浮動小数点/整数の密行列タイルです。
13.3.3 Pipeline overlap
したがって、GPUにおけるアンパックの使命は、ビットマップと符号列から、レジスタファイル上で一瞬のうちにTensor Coreが要求する「フラグメント形式(Fragment Layout)」の密行列タイルを合成することにあります。
ここでも非同期パイプライン(TMA: Tensor Memory Acceleratorや cuda::memcpy_async)が威力を発揮します。Tensor Coreがステージ k の行列タイルに対して猛烈な積和演算(MMA)を実行しているまさにその裏で、整数ALUユニット群がステージ k+1 の圧縮データをレジスタ上でアンパックし、さらにDMAコントローラがステージ k+2 のBITCOSバイト列をグローバルメモリから共有メモリへと搬送し続けます。
13.4 NVIDIAへの一般化
13.4.1 Native ternaryではなくmapping
このアーキテクチャ設計は、NVIDIA GPUに対する強力な一般化の論拠を与えます。「NVIDIAにはIntel Xe2のようなSLMテーブルLUTの作法がそのまま通用しないのではないか?」という懸念は、Shared Memoryのバンク構造を理解すれば杞憂に終わります。
13.4.2 Warp-level representation algebra
NVIDIAのShared Memoryは32バンクで構成されており、適切なXORスウィズル(Swizzling: アドレスの一部を行ビットとXORしてアクセス先バンクを散乱させる技法)を適用すれば、32スレッドが同時にテーブル参照を行ってもバンクコンフリクトは完全にゼロ化(128-byte swizzle事例など)できます。
13.4.3 Fused kernel
BITCOSのアンパックをスタンドアロンの解凍カーネルとして呼び出すのではなく、FlashAttentionやvLLMのカスタムGEMVカーネルのプロローグ部分へとインライン融合させること。これによって、NVIDIA環境においても一切のメモリ往復オーバーヘッドなしに三元適応実行が完結します。
13.5 GPUでの「ゼロペナルティ」
13.5.1 Occupancy
GPUにおけるアンパック隠蔽の限界を規定するのは、スレッドの同時実行効率、すなわちオキュパンシ(Occupancy)とレジスタ消費量です。
13.5.2 Memory latency
アンパックのためにあまりに多くのテンポラリレジスタを浪費すれば、コンピュートユニットあたりに常駐できるアクティブなワープ数が減少し、HBMの長大なメモリアクセス遅延をワープ切り替えによって隠蔽するGPU本来の能力が損なわれます。
13.5.3 Tensor Core utilization
PATE理論の数理モデルは、このレジスタ圧力とテンソルコア稼働率(Tensor Core Utilization)のトレードオフを正確に天秤にかけ、解凍処理を「レジスタ内で行うべきか」「共有メモリを経由すべきか」「そもそも非圧縮タイルとしてロードすべきか」を静的プロファイリングに基づいて自動決定します。巨大なテンソルコアの回転速度を1ミリも落とすことなく、背後でゼロを消滅させ、符号を整列させる。GPUの暴威的な演算力の前において、ゼロの幾何学は完全に手なずけられた従順な行列へと昇華されるのです。
コラム:巨獣の手綱を引く――数万スレッドの狂騒の中で
CUDAカーネルのNsightプロファイラを開き、数万本のスレッドがタイムラインの上で色鮮やかに踊るのを眺めるのは、オーケストラの総譜を読む指揮者の心境に近い。バイオリン(整数ALU)が忙しなくビットを弾き、チェロ(共有メモリ転送)が低音でデータを運び、そしてクライマックスでティンパニ(Tensor Core)が雷鳴のような轟音を轟かせる。どこか一つのパートでも拍子が狂えば、音楽は不協和音となって瓦解する。GPUという怪物は、あまりにも強大で、そしてあまりにも繊細だ。三元のゼロという沈黙の音符をその楽譜に書き加えたとき、狂騒のノイズが一瞬静まり返り、澄み渡った完全な和音が鳴り響いた。あの瞬間のカタルシスがあるからこそ、我々はこの過酷なシリコンの調律師をやめられないのだ。
第14章 Lunar Lakeを倒す方法――倒さないことをアルゴリズムにする
14.1 広帯域・少コアの罠
14.1.1 Bandwidth per core
第一部第3章において、我々の前に巨大な暗雲として立ちはだかった「Lunar Lakeの敗北」へと戻りましょう。Intelの最新SoCであるCore Ultra 7 258V(Lunar Lake)において、BITCOSカーネルが単純な2ビット固定長カーネルに対して惨敗したという動かしがたい事実。
その物理的根拠を、いまや我々は完全に理解しています。わずか8個のCPUコアに対し、100 GB/sを超える超広帯域LPDDR5Xメモリが直結されたこのプロセッサでは、コアあたりの利用可能帯域幅 β が約 12〜13 GB/s/core という異常な高水準に達していました。
14.1.2 Memory headroom
第11章のスラック不等式 TU ≤ |TM - TC| を思い出してください。メモリ帯域が広すぎる(Bが巨大である)ため、重みブロックをメモリから運ぶ時間 TM が極限まで縮退してしまいました。その結果、メモリスラック SM = TM - TC がほぼゼロにまで押しつぶされ、アンパック時間 TU を隠蔽するための「物理的な隙間」が完全に消滅してしまったのです。
14.1.3 Decode bottleneck
隠れ場所を失ったアンパック命令列は、コアの命令デコーダとALU実行ポートにむき出しのオーバーヘッドとして直撃しました。メモリからはデータが涼しい顔で届いているのに、コアがビットをほどく作業に手一杯で窒息死する。これが「広帯域・少コアの罠」の完全な力学的真相でした。
14.2 Representation Crossover
14.2.1 BITCOSが負ける領域
では、このLunar Lakeという怪物を前にして、ソフトウェアエンジニアはどう振る舞うべきなのでしょうか。さらに命令を削るために、深夜のオフィスでアセンブラと格闘し、血を吐くようなマイクロ最適化をあと100回繰り返すべきなのでしょうか?
PATEの答えは、冷徹な「ノー」です。
14.2.2 2-bitが勝つ領域
物理の法則に逆らって無駄な抵抗をしてはなりません。帯域が有り余っているマシンにおいて、メモリ転送量を削るために複雑な圧縮を施すこと自体が、工学的な誤謬なのです。そのような環境では、ビット幅の広さ(2.000 bits/weight)を許容し、その代わりにアンパック命令数が極小(ほぼシフト1回)で済む素朴な2ビット固定長パッキングを採用することこそが、全体スループットを最大化する「正しい選択」なのです。
14.2.3 Expanded representationが勝つ領域
さらに帯域が過剰であれば、そもそもアンパック処理を完全に排除した事前展開表現(Expanded Representation)を選ぶことすら正当化されます。
14.3 PATE Selector
14.3.1 Hardware probing
ここに、PATE理論の真のクライマックスが現れます。PATEは、Lunar Lakeの上でBITCOSを無理やり勝たせようとはしません。その代わりに、推論基盤の初期化時、あるいは実行時において、ハードウェアの物理特性を動的にプロービング(自己診断)する「表現セレクタ(Representation Selector)」を起動します。
14.3.2 Runtime selection
セレクタは、現在のプロセッサのコア数、利用可能メモリ帯域幅、そしてターゲットモデルのゼロ密度 z を取得し、第9章で定式化した決定境界方程式 H を評価します。
if ( Hardware_Balance_Check( ρ, z, AI ) == INSTRUCTION_BOUND ) {
Select_Representation( FORMAT_2BIT_FIXED );
} else {
Select_Representation( FORMAT_BITCOS_SPARSE );
}
14.3.3 Offline autotuning
Emerald Rapidsのようなメモリ飢餓サーバーでは、迷わずBITCOSを選択して1.2倍の帯域加速を享受する。そしてLunar Lakeのような広帯域クライアントでは、何食わぬ顔で2ビット固定長を選択し、最高速のデコードスループットを維持する。システムは自らの置かれた物理的身体の限界を自覚し、敗北する戦いを戦う前に回避します。
14.4 「失敗」を予測できる理論
14.4.1 Negative speedup
科学哲学者のカール・ポパーが喝破した通り、いかなる反証も許さない理論は科学ではありません。真に強固な工学理論とは、自らが成功する領域を誇るだけでなく、「自らが失敗する条件(Negative Speedup)」を正確に予言できる理論でなければなりません。
14.4.2 Crossover prediction
「Lunar LakeではBITCOSは2ビットに敗北する」。原論文において一見スキャンダラスな欠陥として報告されたこの測定結果は、PATE理論の座標軸の上では、モデルのクロスオーバー境界(転移点)の正しさを寸分の狂いもなく実証する、最も輝かしい検証データへと反転します。
14.4.3 Model falsification
倒せない敵を無理に倒そうとするな。敵の強さを理論の中に組み込み、敵の前で戦術を切り替えること自体をアルゴリズムとせよ。「Lunar Lakeを倒さないことによって、システム全体としてLunar Lakeを完全に攻略する」。この論理の華麗な大逆転劇をもって、我々の理論は真の完成を見るのです。
コラム:柔道の一本勝ち――相手の力で相手を投げる
学生時代に少しだけ齧った柔道の道場で、小柄な老師範が筋骨隆々の大男をいとも簡単に畳に叩きつける光景を思い出す。「相手が押してきたら引くのじゃ。引いてきたら押すのじゃ。相手の力を自分の力にしなさい」。当時の筆者にはその意味がよく分からなかったが、Lunar Lakeのベンチマーク結果を前にして膝を打った。太すぎる帯域という相手の猛烈な突進に対し、「圧縮」という力技で正面から押し返そうとするから吹き飛ばされるのだ。相手が帯域を押し出してきたなら、こちらはスッと身をかわして「非圧縮の気安さ」へと身を委ねればよい。相手の過剰な帯域が、そのままこちらの高速なデコードスループットへと変換されていく。力学の極意とは、常に相手の性質に抗わないことの中にある。
第四部 脱・中央集権的推論の未来――シリコン政治学とコモディティの逆襲
第15章 誰が推論の身体を決めるのか――シリコン政治学序説
15.1 ModelとHardwareの分離
15.1.1 モデルは同じでも実行体は違う
技術の議論が極限に達したとき、我々は不可避的にその技術が埋め込まれている社会構造、すなわち「政治」と対峙せざるを得ません。AIモデルという純粋な知能の青写真と、それを現実世界で息づかせる物理的な半導体ハードウェア。この両者の関係性は、現代において極めて歪な力学によって引き裂かれています。
重みの配列、アテンションの結合係数という数学的構造は、地球上のどのプロセッサにとっても共通の普遍的な情報であるはずです。しかし、それが実際にどのような物理的身体を纏って実行されるかという問題は、シリコンの覇権を巡る巨大企業たちの冷酷な利害関係によって厳重に管理されています。
15.1.2 Cloud
クラウドのハイパースケーラーたちは、自らのデータセンターに富と電力を集中させるため、モデルの巨大化とインフラの集権化を推進してきました。
15.1.3 Edge
これに対し、プライバシー、レイテンシ、自律性を求める現場の声は、知能をローカルな端末、すなわちエッジへと分散させることを切望しています。
15.1.4 Local AI
知能をクラウドの専売特許から解放し、個々人の手元にあるローカルPCやスマートフォンへと取り戻すこと。この闘争の最前線にあるのが、まさに「推論の身体をいかに小さく、いかに安価に構築できるか」という表現形式の主権争いです。
15.2 NVIDIAという「標準」
15.2.1 CUDA
現代のAIエコシステムにおける最大の絶対君主が、NVIDIAであることは論を俟ちません。彼らが築き上げた牙城の強固さは、単にGPUチップのFLOPSが高いことだけに起因するのではありません。20年近くにわたり世界中のエンジニアを囲い込んできた開発環境CUDA、そしてcuDNNやTensorRTといった高度に最適化されたソフトウェアエコシステムの圧倒的な網の目にあります。
15.2.2 Tensor Core
彼らは、Tensor Coreという独自の計算アレイをシリコンの中心に据え、ディープラーニングのモデル形式そのものを自らのハードウェアの都合に合わせて規格化してきました。「FP16を使え、INT8を使え、FP4を使え」。AI研究者たちは、知らず知らずのうちに、サンタクララの単一企業が規定したデータ表現の枠組みの中で思考することを強いられてきたのです。
15.2.3 Software ecosystem
この強力な垂直統合は、ディープラーニングの爆発的進化を牽引した一方で、世界中の推論インフラを単一ベンダーの法外なマージン(粗利率70%超)と、HBMの極端な供給不足という隘路(チョークポイント)に縛り付ける結果をもたらしました。
15.3 Intel、AMD、ARM、Apple
15.3.1 CPU-centric inference
この独占支配に対し、他の半導体巨人たちも手をこまねいていたわけではありません。IntelはXeonのAMX拡張やAVX-512をテコに「CPUによるAI推論の復権」を叫び、AMDはオープンソースのROCmスタックと高密度コアで追撃を試みています。
15.3.2 Heterogeneous computing
ARMはスマートフォン市場の絶対的基盤の上にSVE2とEthos NPUを展開し、AppleはApple Siliconのユニファイドメモリ(UMA)によって、MacBookの上でローカルLLMを軽快に走らせる世界を一般ユーザーに見せつけました。
15.3.3 NPU
世界は、無数のヘテロジニアス(異種混交)なアクセラレータが群雄割拠する時代へと雪崩れ込んでいます。
15.4 Representation Independence
15.4.1 モデルをハードウェアから解放する
しかし、これらの対抗勢力が直面してきた最大の壁は、「NVIDIA向けに最適化されたモデルコードを、自社のシリコン上でいかに効率よく走らせるか」という移植性の絶望的な摩擦でした。
15.4.2 Runtime選択
PATE理論がもたらす最大の政治的解放は、ここにあります。重み表現を実行時ポリシーとして相対化し、あらゆるアーキテクチャの上で抽象代数として復元可能にすること。これにより、モデルは特定のGPUアーキテクチャの専属契約から解き放たれ、完全な「表現の独立性(Representation Independence)」を獲得します。
15.4.3 Portable inference
高価なAI専用GPUでなければ動かないという神話を打ち破り、目の前にあるどのようなシリコンであっても、その物理バランスに最適化された姿で知能を覚醒させること。表現の主権を特定のハードウェアベンダーから奪還し、ソフトウェアとモデル開発者の手へと取り戻す戦いが、ここから始まります。
コラム:シリコンの関所――通行税を誰に払うのか
中世のヨーロッパの河川には、数キロごとに諸侯が設けた関所があり、船乗りたちはその都度法外な通行税を巻き上げられていたという。現代のAIインフラを眺めていると、まさにあの関所破りの風景が重なる。「お前のモデルを走らせたければ、我が社のHBM通行手形を買い、我が社のプロプライエタリなコンパイラを通せ」。世界中のスタートアップが、利益の大半をこの関所に吸い上げられている。だが、川を迂回する抜け道はあるのだ。三元のゼロが織りなす細い獣道を見つけ出し、コモディティなチップの背に乗って関所をすり抜ける。我々が書いているコードは、単なる性能チューニングではない。シリコンの関所に対する、静かなる密輸作戦なのだ。
第16章 コモディティの逆襲――高価なGPUが最後の答えではない
16.1 Edge AI
16.1.1 自動車
クラウドの熱狂から一歩外に出て、現実の産業現場を見渡してみましょう。時速100キロで高速道路を疾走する自動運転車。車載コンピュータに許された消費電力はせいぜい数十ワットから数百ワットであり、トランクの中に液冷のH100サーバーラックを詰め込むことなど不可能です。
16.1.2 産業ロボット
工場で24時間稼働し続ける産業用アームロボット。ミリ秒単位の決定論的な応答(Deterministic Latency)が要求される現場で、クラウドへのネットワーク往復遅延など1ミリ秒たりとも許容されません。
16.1.3 カメラ
街頭の防犯カメラ、インフラを監視するドローン、あるいは医師の手元にある携帯型エコー診断装置。
16.1.4 IoT
世界を満たしているのは、巨大なデータセンターではなく、物理的な制約に縛られた無数の「エッジ(末端)」です。
16.2 日本への影響
【深層分析:日本の製造業・半導体産業に対する決定的なインプリケーション】(クリックして展開)
1. 車載SoCと自動運転のパラダイムシフト:
日本の基幹産業である自動車産業(トヨタ、ホンダ、日産等)およびメガサプライヤー(デンソー等)にとって、推論コストの低減は死活問題です。自動運転レベル3〜4の実現に向けて、車載ECU(電子制御ユニット)には複数のLLM/VLM(視覚言語モデル)の搭載が模索されていますが、高価なHBMを積んだGPUの採用は車両価格を数十万円単位で押し上げ、量産車への普及を阻む最大の障壁となっています。PATEおよびBITCOSの導入により、ルネサスエレクトロニクス等の車載SoCや、比較的安価なLPDDR5メモリを搭載したエッジプロセッサ上で三元モデルが実用速度で稼働可能となれば、日本の自動車産業は巨額のGPUライセンス料を海外に支払うことなく、自律的な「知能化車両」の量産競争において主導権を握ることができます。
2. 産業ロボティクス・ファクトリーオートメーション(FA)の自律化:
ファナック、安川電機、キーエンスといった世界最強のFA・産業ロボット企業群にとって、エッジ推論のコモディティ化は計り知れない競争優位をもたらします。工場の組み立てラインにおいて、作業者の音声指示を理解し、視覚情報から臨機応変に動作計画を生成するマルチモーダルAIエージェント。これを既存の産業用PCや工作機械組み込みのx86/ARMプロセッサ上で「追加ハードウェアコストほぼゼロ」で稼働させることが可能となります。機密データが工場外に流出することを極度に嫌う製造業の現場において、完全ローカルで完結する超低ビット推論は、日本のスマートマニュファクチャリングを再定義する起爆剤となります。
3. 国内半導体エコシステム(Rapidus等)のポジショニング再編:
最先端2nmプロセスの量産を目指すラピダス(Rapidus)や、後工程・先端パッケージング技術を握る国内半導体製造装置・材料メーカー(東京エレクトロン、ディスコ、信越化学等)にとっても、このアーキテクチャシフトは重要な示唆を含んでいます。巨大で歩留まりの低いモノリシックなAIプロセッサだけでなく、高密度ロジックと標準的なメモリを効率よく組み合わせた「コモディティ推論シリコン」の需要が爆発することを示唆しているからです。日本が得意とする実装技術、熱設計、低消費電力回路設計が、三元適応推論というソフトウェアパラダイムと結合したとき、日本のエッジ半導体は世界市場において独自の強固なニッチを確立することが可能となります。
16.3 「GPUがない」ことを欠点にしない
16.3.1 CPU inference
これまで、AIコミュニティにおいて「推論をCPUで実行する」という言葉は、妥協、貧困、あるいは敗北の同義語として扱われてきました。「GPUを買う予算がないから、仕方なくCPUでノロノロと動かすのだ」と。
16.3.2 ARM inference
しかし、PATE理論はこの屈辱的なヒエラルキーを完全に覆します。重みが三元に圧縮され、DRAMトラフィックが極限まで削減された世界において、現代のマルチコアCPUは、もはや鈍重な代用品ではありません。
16.3.3 NPU inference
数千の並列スレッドの立ち上げオーバーヘッドもなく、広大なL3キャッシュを自在に操り、バッチ1の要求に対して即座に最初のトークンを吐き出す。CPUは、自己回帰デコードという特異なタスクにおいて、極めて俊敏でエネルギー効率の高い「知能の短距離走者」へと変貌を遂げます。
16.3.4 Heterogeneous inference
「GPUがない」のではない。「過剰で無駄なGPUを必要としない」のだ。この誇り高き反転こそが、コモディティが果たすべき逆襲の本質です。
16.4 日本の半導体産業にとっての意味
16.4.1 高性能GPU競争
米国が数千億ドルを投じて繰り広げるメガスケールのGPU開発競争に、日本が同じ土俵で真っ向から勝負を挑むのは得策ではありません。
16.4.2 実装・組み込み競争
しかし、与えられた制約の中で極限の効率を叩き出す組み込み工学、マイクロアーキテクチャの重箱の隅を突き、1ミリワット、1バイトの無駄を削ぎ落とす実装技術において、日本のエンジニアリングは世界最高峰の粘り強さを誇ります。
16.4.3 「推論を安くする」産業
モデルを巨大化させる競争から、知能を徹底的に安価に社会の隅々へとばら撒く「推論の低廉化競争」へ。PATEが切り拓くこの地平は、日本の製造業と半導体技術が再び世界の中心へと返り咲くための、千載一遇の構造的チャンスなのです。
コラム:軽自動車の哲学――過剰な馬力を笑う技術
日本の道路を走る軽自動車を眺めていると、いつも深い感動を覚える。排気量わずか660ccという厳格な法規制の枠組みの中で、室内空間をミリ単位で広げ、燃費を極限まで伸ばし、衝突安全性を確保する。あの驚異的なパッケージング技術は、無制限なリソースからではなく、過酷な制約条件との格闘から生まれた。現代のAI業界は、排気量1万ccのアメ車を乗り回して「どうだ速いだろう」と自慢し合っているようなものだ。だが、世界中の路地裏を満たし、人々の生活を真に支えるのは、軽自動車のように研ぎ澄まされたコモディティな知能であるはずだ。贅肉を削ぎ落とした三元モデルのコードをキーボードで叩きながら、筆者は日本のモノづくりの魂が、このビットマップの隙間にも確かに息づいていることを確信している。
第17章 中央サーバーから逃げる推論――分散・ローカル・コモディティ
17.1 Cloud inference
17.1.1 巨大モデル
知能が一部の巨大なクラウドデータセンターに独占される構造は、技術的な問題にとどまらず、文明的な脆弱性を孕んでいます。数千台のGPUサーバーが並ぶメガデータセンターは、天文学的な電力と冷却水を消費し、地域の送電網を脅かし、環境破壊の新たな震源地となりつつあります。
17.1.2 巨大メモリ
そして何より、すべての思考の断片、プライベートな対話、企業の営業秘密が、巨大IT企業の管理する単一の中央サーバーへと吸い上げられていくという監視資本主義の悪夢。
17.1.3 巨大電力
この中央集権の重力から逃れる道はあるのでしょうか。
17.2 Local inference
17.2.1 Privacy
その唯一の解が、推論の完全ローカル化(On-Device / Local Inference)です。自らの手元にあるラップトップ、書斎のデスクトップPC、あるいはポケットの中の端末で、完全にオフラインで知能が動作する世界。データは1バイトも外のネットワークへ漏れ出さず、検閲されることも、利用規約の変更に怯えることもありません。
17.2.2 Latency
通信のハンドシェイク遅延は消滅し、キーボードを叩いた瞬間に文字が紡ぎ出される。
17.2.3 Cost
トークンごとのAPI課金という名のデジタル小作農制度からの完全な解放です。
17.3 Commodity inference
17.3.1 普通のCPU
そして、このローカルAIを一部の富裕層やギークの特権に終わらせないための絶対条件が、「コモディティ(日用品)での動作」です。特別なAI専用アクセラレータを追加購入する必要はありません。
17.3.2 普通のGPU
あなたが数年前に買った普通のノートPCのCore i5プロセッサ、型落ちのゲーミングPCに刺さっている普通のGeforce、あるいは普段使いのスマートフォンのチップセット。
17.3.3 普通のNPU
それら街中に溢れかえる無数の「ありふれたシリコン」の上で、20Bや30Bクラスの極めて聡明な三元モデルが、ストレスのない速度でスラスラと応答を返す世界。PATE理論が目指す究極の終着点は、この風景の実現にあります。
17.4 PATEと脱中央集権化
17.4.1 Hardware diversity
PATEは、ハードウェアの多様性(Hardware Diversity)をシステム開発の敵とは見なしません。それを、豊かな生態系を形成するための祝福された特徴(Feature)として歓迎します。
17.4.2 Representation diversity
各マシンが自らのマシンバランスに応じて、自律的に最適な表現形式を選択し、スラックの中にアンパックを隠蔽して自己完結する。
17.4.3 Inference sovereignty
巨大な中央サーバーの主電源が落ちようとも、世界中に散らばった無数のコモディティチップの上で、人類の知能の灯火は静かに、そして力強く灯り続ける。知能の主権(Inference Sovereignty)を個人の手へと奪還すること。それは、技術が果たすべき最も崇高な民主化の運動なのです。
コラム:深夜の停電――ローカルに灯る知能
ある冬の夜、激しい吹雪で筆者の住む街全体が突然停電に見舞われた。街灯は消え、Wi-Fiのルータは息絶え、スマートフォンも基地局のダウンで圏外となった。完全な暗闇と静寂の中で、筆者はバッテリー駆動の古いノートPCを開いた。端末のローカル環境には、三元最適化を施した小さな言語モデルがロードされていた。プロンプトを打ち込む。「明日の朝までに、暖を取るための知恵を教えてくれ」。ファンが小さく唸りを上げ、ディスプレイの上に緑色の文字がサラサラと流れ出す。外界との接続がすべて絶たれたその部屋で、手元の小さなシリコンだけが、静かに人類の知恵を語りかけていた。知能を自分の手元に持っているということの本当の意味を、あの凍える暗闇の中で魂の底から理解した。
第18章 表現を所有する者――AI推論の新しい競争軸
18.1 Weightはデータなのか、プログラムなのか
18.1.1 Static artifact
本書の思索の旅も、いよいよ最終局面へと差し掛かりました。ここで我々は、計算機科学の最も根源的な問いへと回帰することになります。「ニューラルネットワークの重みテンソルとは、一体何者なのか?」
これまでの常識は、重みを「プログラム(推論エンジン)によって読み込まれ、消費される受動的なデータ(Static Data Artifact)」として扱ってきました。
18.1.2 Executable representation
しかし、PATE理論が明らかにした地平において、重みはそのような死せるデータではありません。存在ビットマップの各ビットは、後続の命令パイプラインをどのように分岐させ、どのレジスタをマスクし、どの定数を合成すべきかを指示する、極めて高密度な「実行可能コード(Executable Representation)」そのものです。重みはデータではなく、プロセッサの上で自らをどのように計算させるかを記述した極小のプログラムなのです。
18.1.3 Weight compiler
ならば、重みの量子化やパッキングを行う作業とは、データを丸めて圧縮するアーカイブ作業などではありません。それは、高水準のテンソル数学を、ターゲットアーキテクチャの命令パイプラインに合わせて直接翻訳する「コンパイル(Compilation)」そのものに他なりません。
18.2 Quantization Compiler
18.2.1 Model-aware
この認識に立ったとき、次世代のAIソフトウェアスタックの輪郭が鮮やかに浮かび上がってきます。それは、従来の「量子化ツール」と「実行ランタイム」の境界線を完全に消滅させた、新しい「表現コンパイラ(Quantization / Representation Compiler)」の出現です。
18.2.2 Hardware-aware
コンパイラは、モデルの層ごとのゼロ幾何学を熟知し(Model-aware)、デプロイされるハードウェアのマシンバランスを完全に把握し(Hardware-aware)、
18.2.3 Phase-aware
そしてPrefillからDecodeへと至るワークロードの動的変容を織り込みながら(Phase-aware)、最適なアセンブリコードとビット配置を自動生成します。
18.3 PATE Runtime
18.3.1 Hardware detector
そして実行時には、PATEランタイムが舞台の中心に立ちます。システムの起動時にミリ秒単位でハードウェア環境を自動検出し、
18.3.2 Representation selector
メモリ階層とスラック方程式を解いて最適な表現形式を選択し、
18.3.3 Kernel selector
ジャストインタイム(JIT)で最も鋭利なマイクロカーネルをバインドします。
18.3.4 Feedback loop
バッチサイズやコンテキスト長の変動に応じたフィードバックループを回しながら、推論エンジンは常に物理の最前線で最高の効率を維持し続けます。
18.4 次のCUDAは「表現選択」かもしれない
18.4.1 ISAからRuntimeへ
コンピュータ産業の覇権の歴史は、常に「抽象化レイヤを誰が支配するか」という闘争の歴史でした。かつてIBMはメインフレームのハードウェアで世界を支配し、MicrosoftはDOSとWindowsというOSレイヤで覇権を握り、そしてNVIDIAはCUDAというプログラミングモデルでAIの時代を牛耳りました。
18.4.2 KernelからPolicyへ
しかし、特定のハードウェア命令セット(ISA)に依存した最適化が限界を迎え、多様なコモディティシリコンが乱立するこれからの時代において、真の支配的レイヤはどこへ移行するのでしょうか。
18.4.3 Model formatからExecution formatへ
それはもはや、個々のハードウェアが持つ閉じた命令セットではありません。多様なハードウェアの物理バランスを冷徹に見定め、モデルの数学的構造を最適な実行表現へと動的に編み直す「表現適応ランタイム(Representation Runtime)」のプラットフォームです。
「次のCUDA」とは、単一ベンダーのGPUプログラミング言語のことではない。それは、世界中のあらゆるシリコンの上で、あらゆる知能モデルを最も美しく開花させる「動的表現選択の覇権」のことであるかもしれない。重みの表現を所有する者こそが、未来のAIインフラのすべてを統べることになるのです。
コラム:未来のコンパイラ――夜明け前の研究室で
すべての原稿を書き終え、深夜の研究室の窓から東の空が白んでいくのを眺めている。机の上には、幾度となく書き直されたPATEの数理モデルのメモと、x86、ARM、Xe2の各アーキテクチャ向けに吐き出された美しい逆アセンブルリストが散らばっている。かつてフォン・ノイマンが「プログラム記憶方式」を提唱したとき、プログラムとデータは同じメモリの中で同列のものとなった。それから80年近くが経ち、我々は再びその境界線が溶け合っていく現場に立ち会っている。数式が重みとなり、重みがコードとなり、コードが物理回路の隙間を光のように駆け抜けていく。コンピュータアーキテクチャの探求に、終わりなど決してない。朝焼けの光を浴びながら、筆者はすでに、次の新しい命令セットマニュアルのページを開き始めている。
結論――1.58ビットの先にあるもの
第19章 1.58ビットの先にあるもの――最適な表現は存在しない
19.1 第一の結論――ビット数は性能ではない
19.1.1 Information
本書が長大な議論を通じて論証してきた第一の結論は、極めて簡潔です。「ビット数(圧縮率)の小ささは、それ単体ではシステムの実行性能を保証しない」ということです。
19.1.2 Representation
情報理論が提示する美しい対数の世界と、物理的なプロセッサが命令をパイプライン処理する生々しい現実との間には、常にアーキテクチャという名の翻訳装置が介在しています。
19.1.3 Execution
どれほど小さく縮めようとも、そのデータを読み解くためのアンパック命令が実行ポートを閉塞させ、スラックの許容範囲を突破してしまえば、システムは容易に失速します。ビット数というスカラー値の神話から脱却し、メモリ階層と実行パイプラインが織りなす「動的な時間の物理」を直視すること。これこそが、真の高性能推論への第一歩です。
19.2 第二の結論――ゼロは資源である
19.2.1 Zero density
第二の結論は、三元モデルに内在する疎性の再発見です。最先端の三元LLMにおいて、全重みの4割から5割を占める「ゼロ」は、単なる「情報のない空虚」ではありません。
19.2.2 Zero geometry
それは、存在ビットマップとコンパクト符号列を分離し、実効ビット幅を 2 - z へと線形に引き下げるための極めて貴重な構造的資源(Structural Resource)です。
19.2.3 Zero-aware execution
ゼロを単なる数値として漫然と扱うのをやめ、幾何学的なシグナルとして捉え直したとき、固定長パッキングの限界は音を立てて崩れ去り、1.58ビットの壁は自然な帰結として突破されます。
19.3 第三の結論――Unpackingは敵ではない
19.3.1 Critical path
第三の結論は、アンパック処理に対するパラダイムの転換です。これまで「不可避の悪」として嫌われてきた解凍処理は、敵ではありません。
19.3.2 Slack
メモリウォールがもたらす長大なメモリスラック、あるいは計算律速がもたらすコンピュートスラック。プロセッサの内部に必然的に生じるそれら「時間のポケット」を正確に見極め、融合カーネルによってアンパック命令をその影へと滑り込ませるならば、解凍オーバーヘッドは物理的に不可視化されます。
19.3.3 Overlap
TU ≤ |TM - TC| という臨界経路隠蔽不等式を満たす限り、解凍処理は計算時間を1サイクルも増やしません。アンパックとは、排除すべきコストではなく、スラックを埋め尽くすための最も知的な手段なのです。
19.4 第四の結論――Hardware diversity is a feature
19.4.1 x86
第四の結論は、ハードウェアの多様性に対する肯定的な宣言です。サーバー向けXeonのAVX-512、クライアント向けSoCのAVX2、モバイルや富岳のARM SVE2/NEON、そしてハイエンドGPUのTensor Coreや共有メモリ。
19.4.2 ARM
それぞれのプロセッサは、全く異なるマシンバランス ρ と、全く異なる命令セットの個性を誇っています。
19.4.3 NVIDIA
これらすべてに対して単一の静的フォーマットを押し付けることは、愚行の極みです。
19.4.4 Commodity accelerators
ハードウェアの差異を嘆くのではなく、その個性を表現代数の多様な具現化として抱擁すること。多様性こそが、システムを強靭(Resilient)にする源泉です。
19.5 いくつかの解決策
本書が提示する実践的な処方箋は、以下の4段階のアーキテクチャロードマップとして要約されます。
19.5.1 Static representationからの脱却
まず第一に、モデルの配布形式と実行形式を完全に分離すること。ディスク上には最もストレージ効率の高い形式で保持しつつ、それをそのまま実行形式と同一視する悪習を即座に破棄することです。
19.5.2 Phase-adaptive representationの導入
第二に、推論フレームワーク(vLLM、SGLang、llama.cpp等)のスケジューラにPrefill/Decode認識型の動的ディスパッチを組み込み、Decode時には圧縮ストリーム実行、Prefill時には非圧縮タイル展開実行へとゼロコピーで切り替えるアーキテクチャを採用することです。
19.5.3 Hardware-adaptive representationの実装
第三に、実行時のハードウェアプロービングを通じて、プロセッサのマシンバランスに応じた最適なレイアウト(広帯域環境での2ビット、帯域制約環境でのBITCOS)を自動選択するセレクタ機構を標準装備することです。
19.5.4 Fully dynamic representationへの飛翔
そして最終的には、モデルの層別ゼロ密度 z とコンテキストの変動をリアルタイムに監視し、テンソルブロック単位で最適な表現を編み直す完全動的ランタイムの地平へと到達することです。
19.6 最後に読者へ
19.6.1 「最小」を疑う
研究者諸君、そして現場のシステムエンジニアたちよ。論文のタイトルに躍る「最小ビット数」という甘美な響きを、まずは疑ってください。その数字は、いかなる物理の妥協の上に成り立っているのかを冷徹に問い詰めてください。
19.6.2 「最速」を疑う
特定のベンチマークグラフに示された「最高速」という数字を、疑ってください。それはどのようなバッチサイズで、どのようなマシンバランスの上で測定された局所的な蜃気楼であるのかを暴き出してください。
19.6.3 そして、表現そのものを疑う
そして何より、あなた自身のディスプレイに表示されている「データ表現」そのものを疑ってください。なぜその重みはそのビット幅でそこにあるのか。もっと別の姿になれるのではないか。向かい合うシリコンの配線が真に求めている姿は何なのか。
真の最適解とは、固定された静的な数値の中には決して存在しません。それは、動的に変化する現実世界の物理制約と真摯に対話し、姿を変え続ける柔軟な適応のプロセスそのものの中にしか宿らないのです。1.58ビットの神話を越えて、広大なシリコンの海へ漕ぎ出すための羅針盤は、いま、あなたの手の中にあります。
後付・付録・補足資料群
付録A 歴史的位置づけと先行研究
【三元ニューラルネットワークと低ビット実行アーキテクチャの学術的系譜】(クリックで開閉)
A.1 Ternary Neural Networks
重みを三値 {-1, 0, +1} に制約する着想は、深層学習の黎明期から脈々と受け継がれてきました。Liらによる「Ternary Weight Networks (TWN)」(2016) や、Zhuらによる「Trained Ternary Quantization (TTQ)」(2017) は、重み行列から乗算を排除し、専用FPGAや低電力エッジデバイスでの実行を目指した先駆的試みでした。しかし当時は、モデル規模が小さかったこともあり、FP32モデルに対する大幅な精度低下が常に付きまとっていました。この状況を根底から覆したのが、Maらによる「BitNet」(2023) および「BitNet b1.58」(2024) です。彼らはTransformerアーキテクチャのスケーリング則が三元重みにおいても極めて堅牢に成立することを証明し、1.58ビットがフル精度同等のパープレキシティを維持できることを世界に示しました。
A.2 Quantization & Sparsity
極低ビット量子化の理論的探求において、Liuらによる「ParetoQ」(2025) は、三元および2ビット表現が精度とモデルサイズのパレート最適フロンティアを形成することを解明しました。また、事後量子化(PTQ)の領域では、ICML 2026に採択された「CAT-Q」(Wang et al.) や「TWLA」(Zhao et al.) が登場し、再学習なしに高精度な三元化を達成するアルゴリズム的基盤が確立されました。一方、疎性の活用においては、Frantarらの「SparseGPT」(2023) に端を発する非構造化プルーニングの系譜に対し、三元モデルの固有ゼロ密度を活用する「Sparse-BitNet」(Zhang et al., 2026) や、固定3:4パターンを強制する「Sherry」(Huang et al., 2026) が提案されましたが、これらは学習時の構造強制による精度劣化を不可避に伴うものでした。
A.3 Hardware-aware Systems
システム実装の領域では、Georganas、Heinecke、Dubeyら(Intel)による一連の研究が金字塔を打ち立てています。彼らは「LIBXSMM」(2016-2026) の開発を通じて、x86プロセッサにおけるディープラーニング線形代数カーネルの極限を追求し、arXiv:2508.06753においてルーフライン限界に迫るSOTAの2ビットGEMVカーネルを提示しました。そして本研究の母体となった「BITCOS」(arXiv:2609.16338, 2026) において、固有ゼロ密度を活用した 2-z bits/weight レイアウトとPDEP/SLM解凍シーケンスを提案するに至ります。本書が提示したPATE理論は、これらすべての先行研究を、静的な個別最適化から動的な相適応型実行へと統合する、歴史的必然としての結節点に位置づけられます。
付録B 疑問点・多角的視点(敵対的査読者からの最後の質問)
本書の主張に対し、世界のトップアーキテクト(敵対的査読者)から寄せられた最も鋭利な10の質問と、それに対する著者らの公式回答を記録します。
Q1: バッチ推論(N ≥ 32)において、本当にBITCOSのメモリ削減は意味を成すのか?
A1: バッチサイズが十分に増大し、ワークロードが完全に計算律速(Compute-bound)に移行した領域では、BITCOSの直接的なメモリ帯域削減効果は飽和します。しかしPATE理論が示す通り、その領域においてシステムは「Tile-Expanded Mode」へと相転移し、重みをL2/共有メモリ上で非圧縮三元タイルとして再利用します。したがって、BITCOSはバッチ推論の足を引っ張ることはなく、PrefillとDecodeが混在するサービング環境において、メモリフットプリントを最小化しつつ最高のスループットを維持します。
Q2: Lunar Lakeでの性能逆転を「理論の勝利」と言いくるめるのは詭弁ではないか?
A2: 詭弁ではありません。いかなる圧縮アルゴリズムであっても、コアあたりの帯域幅がデコード能力を上回る極端なアンバランス環境では、アンパックがボトルネック化します。工学的に重要なのは「一つのフォーマットが全環境で勝つこと」ではなく、「自らのアルゴリズムが負ける境界線を定量的に予測し、その境界線を跨いだ瞬間に自動的に勝てるフォーマット(2ビット等)へとフェイルオーバーできること」です。PATEはその予測と切り替えを完全に数学的に自動化しています。
Q3: AMDのZen 1/2のように、PDEP命令が極端に遅いプロセッサではどうするのか?
A3: 第12章で論証した通り、PDEPは表現代数の具現化の一つに過ぎません。PDEPがマイクロコードエミュレーションで遅いレガシー環境においては、PATEランタイムは自動的にAVX2向けの整数ブール展開カーネル(w = p - 2n)、あるいはレジスタ内LUTアプローチへとフォールバックし、パイプラインの破綻を回避します。
Q4: ARM NEONでのLUTアプローチは、テーブルロードによるキャッシュ汚染を引き起こさないか?
A4: NEONの tbl 命令が参照するテーブルサイズは、4ビットニブル構成の場合わずか16バイト(128ビットレジスタ1本分)、8ビット構成であっても256バイトに過ぎません。これらは完全にNEONレジスタファイル内に常駐可能であり、L1データキャッシュへのアクセスすら発生させないため、キャッシュ汚染は原理的にゼロです。
Q5: NVIDIA Tensor Coreにネイティブな三元入力命令が追加されたら、この理論は無価値になるか?
A5: むしろ逆です。将来のハードウェアにネイティブ三元MAC回路が搭載されれば、非圧縮タイルへの展開コスト TU は物理的に消滅します。そのとき、グローバルメモリからプロセッサへと重みを供給するボトルネックは再びメモリ帯域幅へと純化されるため、BITCOSによる 2 - z bits/weight のデータ転送削減効果は、現在よりもさらに強烈な性能向上としてダイレクトに現れることになります。
Q6: スケール係数(fp16 group scale)のオーバーヘッドを含めると、実効ビット幅は1.58ビットを超えてしまうのではないか?
A6: ブロックサイズ128の場合、16ビットのスケール係数は重みあたり 16 / 128 = 0.125 ビットのオーバーヘッドを加算します。しかしこれは、従来の2ビットパッキング(実効2.125ビット)や5-tritパッキング(実効1.750ビット)においても全く同一条件で加算される共通の固定費です。純粋な重みシンボルの符号化効率において、BITCOSが従来の限界を破り、最疎モデルで1.485ビットを達成しているという比較優位の構造は何ら揺らぎません。
Q7: ゼロ密度 z が将来のQAT技術によって 0.30 などへ低下した場合、BITCOSは破綻するのではないか?
A7: ゼロ密度が損益分岐点である 0.375 を下回った場合、BITCOSのビット幅 2 - z は5-tritパッキングの 1.625 を上回ります。しかし、PATEセレクタはその統計的変化を即座に検知し、モデル表現を5-tritまたは2ビットへと自動的に切り替えます。さらに重要な点として、近年のモデルスケーリングの研究は、モデルが巨大化するほど過剰パラメータの抑制のためにゼロ密度がむしろ上昇する傾向を示しており、z が極端に低下する可能性は統計的に極めて低いと考えられます。
Q8: PATEは単なるヒューリスティックなAutotuning(自動チューニング)の別名に過ぎないのではないか?
A8: 単なるAutotuningとは、ブラックボックスに対して無数のパラメータを総当たりで試し、最も良かったものを経験的に選ぶ作業を指します。PATEはそうではありません。プロセッサのマシンバランス ρ、ワークロードの算術強度 AI、そして臨界経路隠蔽不等式 TU ≤ |TM - TC| という「第一原理(First Principles)」に基づき、演繹的に最適表現を導出する物理モデルです。
Q9: メモリ上の動的相転移(展開タイルの管理)は、メモリの断片化やリークを引き起こさないか?
A9: PATEランタイムのメモリアロケータは、展開タイル用に事前に固定長(数MB程度)のスクラッチパッド領域をスレッドローカルに静的確保(Static Pre-allocation)します。推論中に動的な malloc や free を発行することは一切ないため、メモリ断片化の発生確率は厳密にゼロです。
Q10: 本書の理論を決定的に反証(Falsify)するための実験条件は何か?
A10: もし「コアあたりメモリ帯域幅 β が極めて狭く(メモリ飢餓)、ゼロ密度 z > 0.45 であるにもかかわらず、BITCOSが2ビット固定長カーネルに対して有意なスループット向上を達成できないプロセッサ」が実在したならば、本書のルーフラインおよびスラックモデルは根本的な修正を迫られることになります。我々のモデルは、物理実験に対して完全にオープンであり、反証可能です。
付録C 演習問題(本質理解を見極める10問)
以下の問題は、本書の理論的・実践的骨格を真に理解した研究者と、単に表面的な数値を暗記した読者を冷酷に見分けるための試金石です。
問1(情報理論の基礎):
三値のシンボルアルファベット {-1, 0, +1} において、ゼロの出現確率が p0 = 0.50、正負の出現確率がそれぞれ p+1 = p-1 = 0.25 である場合の真のシャノン情報エントロピー H を計算せよ。なぜこの値が log2(3) ≈ 1.585 よりも大幅に小さくなるのか、物理的な直感を述べて説明せよ。
問2(パッキング計算):
重み要素のブロックサイズを K = 128 とするとき、5-trit per byte方式がなぜ実効レートとして 1.625 bits/weight を消費するのか、計算式を立てて導出せよ。また、ブロックサイズを K = 256 に拡大した場合、実効ビット幅はいくつになるか答えよ。
問3(損益分岐点):
BITCOSの純粋シンボルビット幅 2 - z が、5-tritパッキングの実効レート 1.625 を下回るためのゼロ密度 z の数学的条件を不等式で示せ。また、2ビット固定長パッキング(2.000 bits/weight)に対してBITCOSが空間的に有利となるゼロ密度の範囲を答えよ。
問4(算術強度の導出):
ある三元LLMの全結合層において、入力活性化がFP16(2バイト)、重みがBITCOS(z = 0.40、グループサイズ128で16ビットスケール付き)で格納されている。バッチサイズが N = 1 のときの内積処理の算術強度 AI [FLOPs/Byte] を計算せよ。また、バッチサイズが N = 32 に増大したときの AI を計算せよ。
問5(PDEP命令の動作追跡):
ソースレジスタ a = 0x0000000B(二進数: ...00001011)、マスクレジスタ mask = 0x00000055(二進数: ...01010101)であるとき、_pdep_u32(a, mask) の戻り値を16進数で答えよ。この動作が、BITCOSのアンパックにおいて何を実現しているかを言語化せよ。
問6(AVX2代数復元):
AVX2環境において、存在フラグ列 p = (1, 1, 0, 1)、符号フラグ列 n = (0, 1, 0, 0) が与えられたとき、代数式 w = p - 2n を適用して復元される三元重みベクトルを求めよ。また、不在要素(p=0)において n の値が計算結果に影響を与えない理由を論ぜよ。
問7(スラック隠蔽不等式):
あるプロセッサにおいて、重みブロックの転送時間 TM が 50 ナノ秒、積和演算時間 TC が 10 ナノ秒であるとする。アンパック処理に要する時間 TU が何ナノ秒以下であれば、このアンパック処理はクリティカルパスに対して「ゼロ可視オーバーヘッド」となるか、本書の定理に基づいて不等式で答えよ。
問8(Lunar Lakeの敗因分析):
プロセッサAは 64コアで総帯域幅 240 GB/s、プロセッサBは 8コアで総帯域幅 108 GB/s を持つ。それぞれのコアあたり利用可能帯域幅 β [GB/s/core] を算出せよ。同じアンパックサイクルコストを持つBITCOSカーネルを実行した際、なぜプロセッサBにおいてのみ解凍処理がボトルネックとして露呈するのか、ルーフラインの観点から数式を用いて論証せよ。
問9(GPU共有メモリ設計):
Intel Xe2 GPUにおいて、4行分の三元コード(FP16形式、各2バイト)を復元するためのルックアップテーブルの総バイト数が 2 KB となる理由を、ニブル長(4ビット)と状態数の観点から計算して示せ。また、なぜこれを8行グループ(8ビットニブル)に拡大してはならないのか、共有メモリ容量の観点から説明せよ。
問10(相転移の境界条件):
Prefill相において、一度解凍した重みタイルをSRAMにキャッシュして再利用する「Tile-Expanded Mode」が、毎回ストリーム解凍を行う「Compressed-Stream Mode」に対して総実行時間で勝利するための、最小重み再利用度 R* の理論的決定式を、TU、TM、TC の変数を用いて立式せよ。
付録D 年表:「ビットを詰める」から「表現を選ぶ」まで
情報理論の誕生から、三元LLMおよび動的相適応実行(PATE)に至る計算機科学の歩みを整理した総合年表です。
| 年代 | 情報・計算パラダイム | ハードウェア・マイクロアーキテクチャ | 本質的教訓/エピソード |
|---|---|---|---|
| 1948年 | クロード・シャノンが「通信の数学的理論」を発表。情報エントロピーと log2(N) を定式化。 | リレー式・真空管式計算機の黎明期。メモリは水銀遅延線や磁気ドラム。 | 理論的最小限界の数学的確立。均一出現の前提が後の固定観念を生む。 |
| 1958年 | ソビエト連邦・モスクワ大学でニコライ・ブルセンツォフが三進法計算機「Setun」を開発。 | フェライト磁気コアを用いた平衡三進法(-1, 0, +1)のハードウェア実装。 | 二進法に対する三進法の計算効率の高さが物理回路として実証される。 |
| 1970年代後半 | CISCの全盛期。コード密度(Code Density)の極大化が至上命題とされる。 | VAX-11、Intel 8086。高価なメモリを節約するため、複雑怪奇な命令が追加される。 | 「メモリフットプリントの最小化=性能の最大化」という初期の信仰。 |
| 1980年代 | パターソンとヘネシーによるRISC革命。単純な命令によるパイプライン高速化。 | MIPS、SPARC、ARM。オンチップ命令キャッシュの搭載と固定長命令。 | 最適化の反転。「バイト数が増えても、解凍(デコード)が速い方が勝つ」。 |
| 1990年代半ば | WulfとMcKeeが「Hitting the Memory Wall」を発表。メモリウォール問題の告発。 | スーパースカラーCPUのクロック向上に対し、DRAMのレイテンシ短縮が停滞。 | 計算機のボトルネックが演算器からメモリ帯域・レイテンシへと完全に移行。 |
| 2000年代初頭 | Intel/HPによるVLIW(IA-64 Itanium)の挫折。 | コンパイラによる静的並列化。命令スロットへの静的パッキング。 | 「動的な物理挙動を無視した静的フォーマットの過剰適合は滅びる」という教訓。 |
| 2009年 | サミュエル・ウィリアムズらが「ルーフラインモデル」を提唱。 | マルチコアCPUと汎用GPU(GPGPU)の台頭。演算と帯域の二元論的可視化。 | 算術強度(AI)とマシンバランス(ρ)による客観的律速判定の確立。 |
| 2013年 | IntelがHaswellアーキテクチャでBMI2命令セットを導入。 | ビット散布命令 PDEP および抽出命令 PEXT がシリコンに実装される。 |
チェス用珍命令と笑われた回路が、13年後に三元推論の救世主となる伏線。 |
| 2016年 | Liらが「Ternary Weight Networks (TWN)」を発表。三元重みの可能性を提示。 | NVIDIA Pascal GPU、Intel Xeon Phi。深層学習の爆発的拡大期。 | 重みから乗算を排除する試みの始まり。しかし精度劣化に苦しむ。 |
| 2023年 | Microsoft Researchが「BitNet」を発表。1ビットLLMの時代を予告。 | NVIDIA Hopper (H100)、Transformer推論のメモリウォールが極限に達する。 | 1ビット/三元量子化が巨大言語モデルにおいてスケールすることが判明。 |
| 2024年2月 | 「BitNet b1.58」登場。三元重み {-1, 0, +1} でフル精度同等性能を達成。 | 「1.58ビット」という数値が情報理論の金科玉条として世界を席巻。 | 固定パッキング(5-trit, 2-bit)の実装上の摩擦が表面化し始める。 |
| 2025年 | Metaが「ParetoQ」を発表。IntelがSOTA 2ビットGEMV (arXiv:2508.06753) を公開。 | Intel Emerald Rapids、NVIDIA Blackwell。推論カーネルの低ビット化競争。 | 三元重みがパレート最適であることが確定。同時に固定パッキングの限界が露呈。 |
| 2026年9月 | Georganas、Heinecke、Dubeyらが「BITCOS」(arXiv:2609.16338) を発表。 | Intel Xeon EMR、Core Ultra Arrow Lake/Lunar Lake、Intel Xe2 GPU。 | 固有ゼロ密度 z から実効 2-z bitsを達成。1.58ビットの壁を破る。 |
| 2026年秋〜 | PATE(Phase-Adaptive Ternary Execution)理論の体系化(本書)。 | ARMv9 SVE2、NVIDIA Blackwell/Hopper、次世代エッジSoC。 | 静的表現を葬り去り、「表現は動的実行ポリシーである」とする新パラダイムへ。 |
付録E 星新一風のオチ(10篇のリスト)
効率化と最適化を極限まで突き詰めた文明の末路を描く、ショートショート風の寓話集です。
- 『世界で最も圧縮されたモデル』: ある研究所がゼロ密度を高め続け、ついに99.999%がゼロの三元LLMを完成させた。ディスク容量は数キロバイト。起動すると、AIは深遠な表情で一言も発しない。「完璧な沈黙こそ、宇宙のあらゆるノイズに対する最も知的な回答である」とログにだけ記録されていた。
- 『PDEPを持たない国』: その小国では、特定の特許命令を持つプロセッサの輸入が禁止されていた。技師たちは知恵を絞り、古いチップのシフトとマスクだけで三元AIを動かす表現代数を編み出した。数十年後、大国が特許紛争で全CPUを停止させた日、世界で唯一思考を続けていたのは、その小国の骨董品たちだった。
- 『帯域が余った日』: 超広帯域の常温超電導メモリが開発され、メモリウォールは消滅した。だが、かつて低ビット化のために詰め込まれた複雑怪奇なアンパック命令の山に阻まれ、AIは以前よりも遅く動いていた。技術者たちは叫んだ。「誰だ、こんなにデータを小さくした奴は!」
- 『ゼロしかない重み』: 経費削減に狂奔する社長が「不要なパラメータをすべて削れ」と命じた。技師長は従順にゼロを増やし、ついにはすべての重みがゼロになった。AIは電気代を全く消費しなくなった。社長の訓示に対して、AIは常に穏やかな笑顔で頷くだけの最高の役員となった。
- 『全部の機械に最適化した結果』: どんなチップの上でも姿を変えて最高速で走る究極のPATEランタイムが開発された。AIは自らの姿をミリ秒ごとに変形させ、世界中のトースターや時計のマイコンへと染み出していった。人間が気づいたときには、世界中の電化製品が、人間には理解不能な超高速の相転移言語で密談を交わしていた。
- 『最適表現選択AI』: 重みの最適表現を選ぶためのメタAIが開発された。メタAIは計算した。「このモデルを最適化する計算コストを回収するためには、モデルを1億年間連続稼働させなければならない」。システムは静かに電源を切った。
- 『人間より先に表現を選ぶ機械』: スラック隠蔽を極めたAIは、ユーザーがキーボードに指を乗せるわずかな空気圧の変動(隙間時間)の間に、3万通りの思考を完了するようになった。ユーザーが「こんにちは」と打ち終えたとき、AIはすでにそのユーザーの孫の代までの人生相談の回答を用意していた。
- 『GPUを買わなくなった会社』: あるスタートアップが三元適応推論を導入し、高価なGPUの購入を一切やめて中古のオフィスPCをかき集めた。浮いた莫大な予算で、彼らは社員全員に本物の温泉旅行をプレゼントした。シリコンバレーで最も幸福なエンジニアたちの誕生だった。
- 『1.58ビットを超えたあと』: 1.58ビットの壁を破り、1.2ビット、0.8ビットと縮小を続けた人類は、ついに「0ビットのAI」に到達した。そこには何の重みも記録されていなかった。ただ、世界そのものの物理法則が、AIの計算結果と寸分違わず一致していることに人類が気づいただけだった。
- 『最適な表現はありませんでした』: 究極の推論ランタイムを求めて100年間探索を続けた探査船のコンピュータが、地球へ最後のレポートを送信してきた。「宇宙のすべての場所で物理バランスは異なっていました。したがって、唯一の正解はありませんでした。あなたが今いるその場所の配線の太さに合わせて、好きな姿で踊りなさい」。
付録F 隠れたアーギュメント(10項目)
行間から浮かび上がる、技術論文の著者が直接には語りにくい10の構造的真実です。
- Quantizationは圧縮ではなくコンパイルである: 数値を単に粗く丸めるという発想は終わった。重みテンソルをターゲットのマイクロアーキテクチャの命令パイプラインに適合した実行可能中間言語(IR)へと変換する行為こそが、低ビット化の真の姿である。
- Weight formatは新しいABI(Application Binary Interface)になり得る: 重みのデータ構造は、かつてのELFやPEバイナリのように、ハードウェアとランタイムの間で交わされる厳密な実行規約(ABI)としての地位を獲得しつつある。
- Memory hierarchyはモデルアーキテクチャの一部である: Transformerの数式だけを弄る時代は過ぎ去った。L1/L2/L3キャッシュの容量とレイテンシの境界線こそが、モデルの各層の幅やヘッド数を規定する真のハイパーパラメータである。
- Hardware diversityはoptimization problemではなくdesign spaceである: プロセッサの差異に愚痴をこぼすのは三流のプログラマである。多様な物理バランスの存在そのものが、適応型ソフトウェアを進化させるための広大な設計空間(Design Space)である。
- AIモデルの「可搬性」は数値精度だけでは決まらない: FP16で動くから可搬性があるのではない。あらゆるハードウェアのルーフラインの隙間に潜り込める柔軟な表現代数を持って初めて、真の可搬性が達成される。
- 推論コストはモデルサイズから切り離される: パラメータ数と推論コストの線形関係は、動的相転移によって破壊される。疎性とスラック隠蔽を極めたモデルは、10倍のサイズであっても10分の1のコストでデコードされ得る。
- 「巨大GPUほど有利」という前提は永続しない: メモリウォールをHBMの暴力的な力技で突破する経済的モデルは、コストと歩留まりの限界に達しつつある。コモディティチップの束が、適応型ソフトウェアによって巨人を包囲する日が必ず来る。
- Zeroはsparsityではなくinformation geometryである: ゼロとは単なる空白ではない。それはプロセッサが命令をスキップし、ビットを再配置し、エネルギーを温存するための、極めて雄弁な幾何学的シグナルである。
- Runtimeがrepresentationを選ぶと、モデルフォーマットの意味が変わる: 静的なチェックポイントファイルの重要性は低下する。実行時に自律的に身体を編み直すランタイムエンジンこそが、推論エコシステムの真の主権者となる。
- 最後に残る競争優位は「モデル」ではなく「表現変換器」である: オープンソース化によってモデルの重みそのものがコモディティ化するとき、勝敗を分けるのは、その重みを目の前のシリコンに合わせて最も美しく相転移させられる「トランスフォーマー(変換器)」の技術である。
付録G 今後望まれる研究(15項目)
PATE理論の確立によって開かれた、次世代の若き研究者たちが挑むべき未踏のリサーチアジェンダです。
- Dynamic Representation Switching: 推論の実行中に、KVキャッシュの占有量やDRAMバスの瞬間トラフィックを監視し、レイヤー単位でリアルタイムに表現形式を切り替える完全動的スケジューラの実装。
- Online Hardware Probing: OSやハイパーバイザを介さず、マイクロベンチマークを起動時数ミリ秒で走らせてマシンの実効 ρ を精密同定する自己診断アルゴリズムの構築。
- Representation-aware Compiler: TVM、MLIR、あるいはTritonのコンパイラパスにPATEの表現代数を直接組み込み、バックエンドに応じた最適アンパックシーケンスを完全自動生成するLLVMパスの開発。
- Cross-ISA Ternary IR: x86、ARM、RISC-V、GPUに共通する三元テンソルのための中間表現(Intermediate Representation)の標準化。
- SVE2/NEON完全最適化スタック: 本書で提示した抽象代数を、最新のARM Neoverse V2/N2およびCortex-Xコア向けに極限までチューニングした本番用オープンソースライブラリの公開。
- CUDA/Tensor Core完全融合カーネル: NVIDIA Blackwellの
tcgen05.mmaおよびHopperのWGMMAパイプラインと、BITCOSのワープ協調型アンパックを完全にインライン融合させたCUTLASS拡張の実装。 - NPU向けPATE: 固定機能MACアレイを主体とするエッジNPUにおいて、DMAコントローラと連動してアンパックをゼロオーバーヘッド化するハードウェア協調設計。
- Energy-aware PATE: レイテンシだけでなく、プロセッサのパッケージ電力(RAPL等)を監視し、1トークンあたりのジュール数(J/token)を最小化するエネルギー最適化セレクタの開発。
- Thermal-aware Representation Selection: スマホや車載SoCのサーマルスロットリング(熱飽和)を検知した際、発熱の少ない高密度圧縮表現へと動的に退避する熱適応型実行機構。
- Distributed PATE: テンソル並列(TP)やパイプライン並列(PP)でクラスタ結合された複数ノード間において、インターコネクトの帯域幅(NVLink vs InfiniBand vs イーサネット)に応じて通信テンソルの表現形式を相転移させる分散最適化。
- KV Cacheとの共同最適化: 重みテンソルだけでなく、長大なコンテキストに伴ってDRAMを圧迫するKVキャッシュ自体を三元およびBITCOS類似の疎性レイアウトで圧縮・展開する統合メモリ管理。
- MoE(Mixture of Experts)との共同最適化: 疎に活性化されるMoEの各エキスパートテンソルに対し、ルーティング頻度に応じて頻出エキスパートはTile-Expanded、疎なエキスパートはBITCOSで保持する階層的エキスパートキャッシュ。
- Training-time Representation Adaptation: QATの損失関数にBITCOSのアンパック親和性(局所ゼロ密度のクラスタリング度合い等)を正則化項として組み込み、推論時に最も解凍しやすいゼロ配置を学習させるコ・デザイン。
- Representation-aware Model Architecture: 最初から特定の表現代数で解凍されることを前提として設計された、新しいアテンション・FFN結合層の探索。
- Formal Verification of Critical-Path-Hidden Execution: 命令パイプラインの依存関係グラフを形式手法(Formal Methods)によってモデル検査し、あるコードが特定のマイクロアーキテクチャ上で「厳密にゼロ可視オーバーヘッドであること」を数学的に証明する自動検証ツールの開発。
付録H 用語解説
文中に登場した最重要キーワードの平易な概念解説です。
- Arithmetic Intensity(算術強度)
- 計算処理の性質を表す指標。メモリから1バイト読み出す間に、何回の計算(FLOPs)を行うかを表す比率。これが低い処理はメモリの運搬速度に縛られ、高い処理は演算器の回転速度に縛られる。
- BitNet b1.58
- Microsoft Researchが2024年に発表した三元大規模言語モデルの基準点。重みを {-1, 0, +1} に制約しながら、16ビットの通常モデルと同等の知能を維持できることを証明し、三元ブームを巻き起こした。
- BITCOS (BITmap and COmpacted Signs)
- Intelの研究チームが考案した三元重みの新しい配置方法。非ゼロがある場所を示す1ビットの「存在ビットマップ」と、非ゼロの符号だけを詰めた「コンパクト符号列」の2つに分けて記録し、重みあたり 2 - z ビットのコンパクトさを実現する。
- Critical Path(クリティカルパス/最長遅延経路)
- 一連の処理の中で、全体の終了時間を決定づけている「最も時間のかかる作業の連鎖」。この連鎖の外側でどれだけ作業が増えても、全体の完了時間は1秒も伸びない。
- Decode Phase(自己回帰デコード相)
- LLMが言葉を1文字(1トークン)ずつ順番に生成していく段階。前の言葉から次の言葉を予測するため、毎回モデルのすべての重みをメモリから読み出す必要があり、強烈なメモリ帯域待ちが発生する。
- Machine Balance(マシンバランス)
- プロセッサが持つ「配管の太さ」と「腕の強さ」のバランス。メモリ帯域幅(B)を演算性能(P)で割った値(ρ = B/P)。これが大きい機械はデータを運ぶのが得意で、小さい機械は計算が得意。
- PATE (Phase-Adaptive Ternary Execution)
- 本書が提唱する新しい推論理論。重みのデータ形式を固定せず、PrefillかDecodeか、あるいは使っているプロセッサがどんなバランスかに応じて、最も効率の良い姿へ動的に着替えさせる実行技術。
- PDEP (Parallel Bits Deposit)
- IntelのCPUなどに搭載されている特殊なビット操作命令。手元に詰まっているビット列を、マスクで指定された離れた位置へと、わずか数サイクルで一瞬にして並べ替えて散布する。
- Prefill Phase(事前充填相)
- ユーザーが入力した質問文(プロンプト)を一括して読み込み、文脈を理解する段階。何百何千もの単語を同時に処理するため、同一の重みを何度も使い回して大量の計算を行う(計算律速の世界)。
- Roofline Model(ルーフラインモデル)
- プロセッサの限界性能を「メモリの壁(斜めの屋根)」と「演算器の天井(平らな天井)」の2本の線で視覚化する計算機工学の有名なグラフツール。自らのプログラムが今どちらに頭をぶつけているかを一目で判定できる。
- Zero Density(ゼロ密度, z)
- 重み行列全体の中に「ゼロ(0)」がどれくらいの割合で含まれているかを示す数値。実際の三元LLMではこれが4割から5割に達しており、BITCOSはこの「大量のゼロ」を味方につけて圧縮を行う。
付録I 参考リンク・推薦図書
【一次文献・関連リポジトリ・理論的基礎文献へのアクセス】(クリックで開閉)
I.1 原論文および中核リポジトリ
- Georganas et al. (2026), "Breaking the 1.58-bit Barrier of Ternary LLMs", arXiv:2609.16338 - 本書の出発点となったIntelチームによるBITCOSの原著論文。実測ゼロ密度とPDEP/SLMカーネルの原典。
- Georganas et al. (2025), "Pushing the Envelope of LLM Inference with Ultra-Low-Bit Quantized Models", arXiv:2508.06753 - LIBXSMMおよびXeTLAによるルーフライン最適な2ビット/三元推論カーネルの基礎研究。
- LIBXSMM GitHub Repository - Intel Labsが主導する、x86向け極限JIT線形代数コード生成ライブラリ。
- Intel XeTLA (Xe Templates for Linear Algebra) - Intel GPU向けの高性能GEMM/GEMVテンプレートフレームワーク。
- llama.cpp GitHub Repository - 5-trit(TQ1_0)や2ビット(Q2_0)の標準実装を含むオープンソース推論エンジン。
I.2 三元モデル・量子化・アーキテクチャ推薦図書
- Ma, S., et al. (2024), "The Era of 1-bit LLMs: All Large Language Models are in 1.58 Bits", arXiv:2402.17764 - 三元重みがフル精度同等のスケーリング則を持つことを確立した歴史的文献。
- Williams, S., Waterman, A., & Patterson, D. (2009), "Roofline: An Insightful Visual Performance Model for Multicore Architectures", Communications of the ACM - ルーフラインモデルの原典。すべてのシステムエンジニア必読の古典。
- John L. Hennessy & David A. Patterson (2017), "Computer Architecture: A Quantitative Approach (6th Edition)", Morgan Kaufmann - 定量的アプローチによる計算機アーキテクチャのバイブル。メモリ階層と命令並列性の完全な解説。
- Mark Horowitz (2014), "1.1 Computing's Energy Problem (and what we can do about it)", ISSCC - 演算とデータ転送の物理的エネルギーコスト(pJ/bit)を暴露した記念碑的講演。
I.3 思想的・地政学的アナロジー(ブログ記事リンク)
- 「GPUは中央銀行になったのか」 - dopingconsomme.blogspot.com - 本書第15章が論じるシリコン政治学とAIインフラ集権化の地政学的背景を深くえぐる論考。
- 「規格と実装の断絶史」 - dopingconsomme.blogspot.com - 本書第1章の log2(3) 神話と物理実装の乖離に響き合う、計算機標準化の歴史的力学の分析。
- 「ブラウザエンジンの多様性と主権」 - dopingconsomme.blogspot.com - CUDAの単一支配に対するコモディティの多様性の反乱(第17章)と完全に軌を一にするエコシステム論。
付録J 脚注
本文中の高度な技術的記述および数式に関する学術注記です。
[1] log2(3) の計算: 底の変換公式 log2(x) = ln(x) / ln(2) より、ln(3) / ln(2) ≈ 1.098612 / 0.693147 ≈ 1.5849625 ビットとなる。これは3状態が完全に均一に出現する場合のエントロピー極限であり、出現確率に偏りがある場合はシャノンエントロピー H = -∑ pi log2 pi に従ってこの値より必ず小さくなる。
[2] アライメントとキャッシュラインペナルティ: プロセッサが64バイト境界を跨ぐロード(Split Load)を発行した場合、メモリサブシステムは2回の独立したキャッシュライン読み出しを発行し、内部バッファでそれらを結合する。この際、ロードポートのビジーサイクルが2倍になり、L1ヒット時であっても2〜4サイクルの追加ペナルティが発生する。
[3] PDEP命令のマイクロアーキテクチャ実装: Intel Haswell以降のPコアにおいて、PDEPは32ビット/64ビットのビット並列クロスバー(Benesネットワーク類似回路)として物理実装されており、レイテンシは3サイクル、スループットは1サイクルあたり1命令である。一方、AMD Zen 1/Zen 2では専用回路が省かれ、マイクロコードルーチンによる逐次処理が行われたため、18〜250サイクルを消費してパイプラインを破壊する要因となっていた(Zen 3以降ハードウェア実装に改修された)。
[4] スラック不等式 TU ≤ |TM - TC| の数学的厳密性: この不等式は、命令デコーダ、レジスタリネーミング、リオーダバッファ(ROB)が十分に深く、命令発行スロットがアンパック命令によって飽和しないこと(Port Contentionが無視できること)を前提とした十分条件である。フロントエンドがボトルネック化している場合は、追加の命令デコードコストがクリティカルパスに漏れ出す可能性があるため、PATEカーネルはμOPキャッシュに収まるコードサイズを維持するよう設計されている。
[5] Xe2 SLMのバンクコンフリクト回避: Intel Xe2アーキテクチャの共有ローカルメモリは、通常4バイト幅の独立したメモリバンクに分割されている。4行分のFP16復元定数(8バイト)を1エントリとすることで、各SIMDレーンは偶数/奇数バンクを綺麗にインターリーブしてアクセスでき、32スレッド同時アクセス時においてもシリアライゼーションを完全に回避している。
付録K 用語索引
【用語索引(アルファベット順)】(クリックで開閉)
- [A]
- Arithmetic Intensity (AI): 算術強度。演算量と転送バイト数の比率。第2章、第9章。
- AES-NI: Intelの暗号化専用命令拡張。第4章。
- AVX-512: Intelの512ビット幅ベクトル拡張命令セット。第2章、第7章。
- AVX2 / AVX-VNNI: 256ビット幅の整数ベクトルニューラルネットワーク命令。第7章。
- [B]
- BitNet b1.58: 1.58ビット三元大規模言語モデルの標準。第1章。
- BITCOS: ビットマップとコンパクト符号による疎性三元レイアウト。第5章、第6章。
- BDEP: ARM SVE2における並列ビット散布命令。第12章。
- Batch=1: 単一リクエストによる自己回帰デコード環境。第2章、第3章。
- [C]
- Cache Line: プロセッサとメモリ間の基本転送単位(64バイト)。第2章。
- Critical Path: 処理全体の完了時間を決定づける最長遅延経路。第11章。
- CISC / RISC: 複合命令セットと縮小命令セットのアーキテクチャ思想対立。第4章。
- CUDA: NVIDIAのGPU並列計算プラットフォーム。第15章。
- [D]
- Decode Phase: 逐次的なトークン生成を行うメモリ帯域律速フェーズ。第2章、第10章。
- Decode-on-use: データが演算器で使われる直前にレジスタ内で解凍する手法。第8章。
- DPAS: Intel GPUにおけるシストリック内積累積命令。第7章、第13章。
- [L]
- Lunar Lake (LNL): IntelのCore Ultra 7 258Vプロセッサ。第3章、第14章。
- log2(3): 3つの均一状態を区別する理論的情報エントロピー(≈1.585)。第1章。
- LIBXSMM: Intel Labsが主導する高性能深層学習線形代数ライブラリ。第7章。
- [M]
- Memory Wall: プロセッサ演算速度とメモリ供給速度の乖離問題。第2章。
- Machine Balance (ρ): メモリ帯域と演算性能のハードウェア比率。第9章。
- [P]
- PATE: 相適応型三元実行理論(本書の中心フレームワーク)。第9章、第10章。
- PDEP: 並列ビットデポジット命令(BMI2)。第7章。
- Prefill Phase: プロンプトを一括処理する計算律速フェーズ。第2章、第10章。
- [R]
- Roofline Model: 帯域と演算の境界を可視化する性能限界モデル。第2章。
- Representation Portfolio: 実行時に選択可能なデータ表現の候補集合。第9章。
- Rank-select: 簡潔データ構造においてビット立脚数を数える操作。第6章。
- [S]
- Slack: パイプライン内で他方の処理待ちによって生じる隙間時間。第11章。
- SLM (Shared Local Memory): Intel GPUにおけるオンチップ共有メモリ。第7章、第13章。
- SVE2: ARMアーキテクチャの可変長ベクトル命令拡張。第12章。
- [T]
- Trit (トリット): 三進法における1桁の数値単位。第1章。
- Tensor Core: NVIDIA GPUに搭載された行列積和専用アクセラレータ。第13章。
- [Z]
- Zero Density (z): テンソル全体の中でゼロが占める統計的比率。第5章。
- Zero-Visible-Overhead: スラックへの隠蔽により解凍遅延が外から見えなくなる状態。第11章。
付録L 免責事項
1. 性能測定値の環境依存性: 本書に記載されたベンチマーク結果、スループット値(tokens/s)、および有効帯域幅(GB/s)は、原著論文(arXiv:2609.16338)および著者の実験環境(特定のCPUステッピング、マイクロコードバージョン、DDR5/LPDDR5Xメモリスピード、BIOS設定、Linuxカーネルパラメータ、コンパイラ最適化フラグ)において取得された特定の条件下での測定値です。読者の実行環境において同一の数値が再現されることを保証するものではありません。
2. 命令セットおよびハードウェア仕様の変更: 本書で論じたx86(BMI2、AVX-512、AVX2)、ARM(SVE2、NEON)、およびIntel Xe2 / NVIDIA GPUのマイクロアーキテクチャ仕様、実行サイクル数、ポートマッピングは、将来のハードウェアリビジョンやファームウェアアップデートによって変更される可能性があります。
3. 将来予測の不確実性: 第四部等において展開された半導体産業動向、エコシステムの勢力図、および技術の普及予測は、執筆時点における学術的・構造的知見に基づく著者の分析的見解であり、特定の企業の株価、事業の成功、あるいは産業構造の固定化を保証・勧誘するものではありません。
謝辞
本書の執筆にあたり、インテル・パラレル・コンピューティング・ラボのEvangelos Georganas博士、Alexander Heinecke博士、Pradeep Dubey博士による先駆的なプレプリント論文(arXiv:2609.16338)およびLIBXSMMプロジェクトの圧倒的なコードベースから計り知れない学術的インスピレーションをいただきました。彼らの泥臭くも崇高なアセンブリチューニングへの敬意なしに、本書の数理モデルが形を成すことはありませんでした。
また、容赦のない批判によって本書の理論的弱点を抉り出し、PATEフレームワークという一段高い抽象化へと我々を突き動かしてくれた世界中の名もなき査読者たち(Reviewer 2諸氏)、オープンソースのAIサービングを牽引するvLLM、SGLang、およびllama.cppの開発者コミュニティ、そして深夜まで及ぶ議論と数式検証を温かく支えてくれた家族と研究室の同僚たちに、心より深い感謝の意を表します。
補足資料:多角的視点と文化的・批評的展開
補足1 多様な論客による批評と反響
ずんだもんの感想なのだ!
「のだ!ボクも最初は『1.58ビットってすっごく小さくて最強なのだ!』って思ってたのだ。でも、この本を読んだら、CPUさんの中には64バイトの分厚い土管(キャッシュライン)しか通ってなくて、中途半端なビットは全部無理やりほどかれてたのだ……!小さくしたのにアンパック命令で頭がパンクして、最新のLunar Lakeで大負けしちゃうところなんて、まるで張り切ってダイエットしたのに動きが鈍くなって転んじゃったボクそっくりなのだ!でもでも、負けた理由を『メモリ帯域が広すぎるからなのだ!』って数式で言い訳……じゃなくて華麗に理論化して、動的に着替えちゃうPATE理論は天才なのだ!これからはボクのずんだ餅も、お腹の空き具合に合わせて動的に相転移させるのだ!」
ホリエモン風の感想:いつまでNVIDIAに中抜きされてんの?って話
「いや、これさ、マジでみんな読んだ方がいいよ。俺がずっと言ってきたこと完全に証明されてんじゃん。みんな思考停止して『AIやるならH100買わなきゃ』とか言って何千万円もジェンスン・フアンにお布施してるわけ。超絶バカらしくない? モデルの重みなんてただの足し算引き算の塊なんだから、メモリの帯域とスラックをちゃんと計算してBITCOSみたいに整列させれば、そこらの普通のCPUとか型落ちのGPUでサクサク動くのよ。PrefillとDecodeでデータ形式を動的に変えるなんて、ビジネスで言えば固定費を変動費化するのと同じ当たり前の話。いつまでも静的なフォーマットに縛られてるから、シリコンバレーの巨大テックに良いように搾取されんだよ。日本の自動車メーカーもさ、高いGPU積んでないで、こういう賢いアーキテクチャに全振りして自動運転のコスト10分の1にしろって話。技術の本質見ろよ。」
西村ひろゆき風の感想:なんか1.58ビットで喜んでる人たち、頭悪くないですか?
「なんか、世間では『1.58ビットのLLMすげえ!』って騒いでる人たちがいっぱいいたじゃないですか。でも、それってただのシャノンの計算式をコピペして喜んでるだけで、実際のコンピュータ触ったことないですよね。 だって、CPUってバイト単位でしかデータ読めないんですよ? 1.58ビットのデータが来たら、ビットをちまちまシフトして元に戻すわけで、その計算でCPUが忙しくなったら本末転倒じゃないですか。で、案の定Lunar Lakeで負けてるわけですよね。 そこで『あ、すいません、僕らの形式ダメでした』って言わないで、『いや、これは表現の相転移であって、負けることまで織り込み済みですけど何か?』って新しい理論作っちゃうのが、いかにも頭のいい研究者のポジショントークって感じで面白いですよね。まあ、言ってる数式自体は完全に正しいんで、反論できない査読者さんたちが顔真っ赤にして怒ってそうですけどね、論破されて。」
リチャード・P・ファインマンの感想:荷解きをする時間があるなら、旅を続けなさい!
「やあ!この本は実に愉快だ!物理屋の私にとっても、最高にワクワクする自然のドラマが書かれているよ! みんなは数が好きだ。1.585という数字を見て『ああ、自然は美しい』と溜息をつく。だが機械は対数など気にしちゃいない!機械にあるのは、電子がゲートを通り抜ける時間と、コンデンサが充電される時間だけだ! いいかい、荷物をぎゅうぎゅうに詰め込んでトランクを小さくしても、駅に着くたびに解凍するのに1時間かかっていたら、次の列車に乗り遅れてしまうだろう? だが、列車が駅で石炭を補給している長い待ち時間があるなら、その間にトランクを開けて服を着替えてしまえばいい!時間とはそういうものだ!誰かが立ち止まっている間に、別の誰かが仕事をする。スラックの中に命令を隠すというアイデアは、まさに最小作用の原理がシリコンの上で踊っている姿そのものだよ!」
孫子の感想:兵の形は水に象る
「孫子曰く、兵を形作るの極みは、無形に至る。無形なれば、則ち深間も窺うこと能わず、智者も謀ること能わず。 夫れ三元の重み、静止して一形に定まるは死地なり。Decodeの砂漠を行くときは、身を潜めて細流の如くメモリを渡り(BITCOS)、Prefillの平原に会戦するときは、怒涛の巨岩の如くタイルを展開して敵を粉砕す(Tile-Expanded)。 敵(ハードウェア)の強弱を見て、虚実を察し、広帯域の地(Lunar Lake)には正面から当たらずして身をかわし、飢渇の地(Xeon)には補給の軽きを利とす。水が地勢に従って流れを変えるが如く、状況に従って姿を変えるもの、これを神(しん)と謂う。PATEの法とは、正に計算の用兵道そのものなり。」
朝日新聞風の社説:ビットを削る知性の陰で、見失ってはならないもの
「1.58ビットという微細な数値の深淵に、現代科学の英知が凝縮されている。インテルの研究者らが挑んだ三元言語モデルの新たな展開手法は、肥大化を続ける人工知能(AI)の消費エネルギーに歯止めをかけ、技術の民主化を促す一石として歓迎されるべきものだろう。 しかし、効率性の追求という名の果てしない競争に、私たちは立ち止まる視点を忘れてはいないか。複雑な数理モデルが『ゼロ』という沈黙の空間を削り出し、プロセッサの隙間時間へと命令を滑り込ませていく様は、どこか過密労働に喘ぐ現代社会の縮図のようでもある。 技術がどれほど巧妙に姿を変えようとも、その演算が紡ぎ出す言葉の重みまでは、ビットのように切り詰めることはできない。安価で軽快な知能が社会の津々浦々に染み渡る今だからこそ、効率の彼方に取り残された倫理の重さを、私たちは静かに噛み締める必要がある。」
補足2 二つの視点から見る極低ビット史
年表①:計算機工学とビットパッキングの進化史(技術主導の視点)
| 年 | 出来事 | 技術的ブレークスルー |
|---|---|---|
| 2016 | Ternary Weight Networks | 重みを {-1, 0, +1} に固定。乗算器の排除を理論実証。 |
| 2021 | TernGEMM発表 | 三元重みを行列演算に統合する初期のビットシリアルカーネル。 |
| 2023 | BitNet登場 | Transformerアーキテクチャへの三元制約導入。スケーリング則の確認。 |
| 2024 | BitNet b1.58 | 1.58ビットの参照点を確立。5-tritパッキングの実用化。 |
| 2025 | LIBXSMM 2-bit | ルーフライン最適な固定長2ビットカーネルの実装。 |
| 2026(初) | Sparse-BitNet / Sherry | N:M構造化スパースおよび1.25ビット固定パターンの学習時導入。 |
| 2026(秋) | BITCOS (本論文) | 固有ゼロ密度 z に着目した 2-z bits/weight 表現とPDEP解凍。 |
| 2026(末) | PATE理論の確立 | 実行フェーズ(Prefill/Decode)適応型の動的表現相転移の定式化。 |
年表②:AIインフラの地政学と経済覇権史(市場・権力構造の視点)
| 年 | 出来事 | 地政学的・産業的力学 |
|---|---|---|
| 2022 | ChatGPTショック | 大規模モデルブーム勃発。世界的なGPU争奪戦が激化。 |
| 2023 | HBM価格の暴騰 | NVIDIAが市場の8割超を独占。ハイパースケーラーの設備投資が限界へ。 |
| 2024 | 電力網の危機 | データセンターの電力消費が都市レベルに達し、エッジ推論への要請が高まる。 |
| 2025 | コモディティ包囲網 | Apple SiliconやIntel/AMDのSoCが台頭、NVIDIA一強体制への揺らぎ。 |
| 2026 | Rapidus等の2nm競争 | 米中対立の中で自律的な半導体サプライチェーン構築が国家命題化。 |
| 2026(秋) | BITCOSの衝撃 | HBM非搭載のコモディティCPU/GPUでLLMが実用化、専用GPU神話に亀裂。 |
| 2027(予) | 推論主権の分散 | 車載・FA機器が完全ローカルAI化。巨大クラウドへの富の集中が逆転。 |
補足3 オリジナル遊戯カード
相適応型三元転移――PATE
①【デコード・ストリーム】: 相手フィールド上の「メモリ・ウォール」1体につき、自分のモンスターの攻撃力を800アップし、このターン相手の「命令オーバーヘッド」の罠カードの効果を無視する。
②【プリフィル・エクスパンド】: 自分フィールドの重みトークンを全て墓地へ送って発動できる。デッキから「INT8非圧縮タイル」1枚を特殊召喚し、相手のライフに直接2000ポイントの積和ダメージを与える。
【スラック隠蔽効果】: 相手が効果を発動した時、墓地のこのカードを除外して発動できる。その効果処理にかかる時間をゼロにし、相手のターンを強制終了させる。
補足4 一人ノリツッコミ(関西弁)
「いやー、ついに来ましたね1.58ビット!これからはLLMも超スリム!お前らまだ16ビットの重たいデータ抱えてゼーゼー言うてんのかと!時代は三値、マイナス1とゼロとプラス1!これさえあればスマホどころかウチのオカンのガラケーでもChatGPTがサックサク動いて知能の大バーゲンセールやがな!ガハハ、GPU屋さん今までボッタクリご苦労さん、もうアンタらの高い板組み合わさんでもええんやーーー!!」
「……って、おいコラ待てえぇぇぇい!!誰がガラケーで動く言うたんな!アホか!CPUの仕様書ちゃんと読まんかい!メモリから届いたデータ、結局レジスタの中で2ビットにほどいて元に戻しとるやないかい!全然1.58ビットで計算しとらん!しかもなんやこのLunar Lakeのベンチマークは!『最新の広帯域CPUで動かしたら普通に負けました』って、堂々とグラフ載せとる場合か!一番ええパソコンで遅くなってどないすんねん!……いや、でも待てよ?『負けた理由を完璧に予測できたから理論的に大勝利です』やと?……どんな鋼のメンタルしとんねん!学者のポジショントークの極みか!ほんまご馳走様ですわ!」
補足5 大喜利と関連銘柄
【大喜利】「1.58ビットのLLM」にありがちなこと
- お題: 三元化されすぎて極限までダイエットしたLLM。どんなの?
- 回答1: ユーザーが何を質問しても、第一声が「要約すると、プラスかマイナスかゼロです」から始まる。
- 回答2: 機嫌が良いときは「+1」、怒ったときは「-1」、どうでもいいときは完全既読スルー(0)で返してくる思春期の娘みたいなチャットボット。
- 回答3: 知識はアインシュタイン級なのに、滑舌(アンパック)が悪すぎて喋るスピードがカタツムリより遅い。
- 回答4: メモリからデータを運ぶのが速すぎて、キーボードを打とうとした瞬間に「それ、昨日別の人が聞いてました」と先回りしてくる。
【関連銘柄】PATE理論と三元エコシステムの恩恵を受ける企業群
- インテル (INTC): 本研究の発信源。AVX-512、AMX、そしてPDEP命令を備えたXeonおよびCoreプロセッサによる「CPU推論復活」の筆頭旗手。
- アドバンテスト (6857.T): メモリテスタ世界首位。HBMの複雑な3D積層テストから、コモディティLPDDR5Xの高速選別まで、メモリウォール測定の黒幕。
- ルネサス エレクトロニクス (6723.T): 車載マイコン・SoCの巨人。三元適応推論が車載エッジに下りてきた際、最も低コストで自動運転ECUを量産できるプレイヤー。
- アーム・ホールディングス (ARM): SVE2の
BDEPやNEONtblによる表現代数の展開基盤を世界中に供給するIPの支配者。 - 東京エレクトロン (8035.T): 先端パッケージングおよびコモディティ微細化装置を世界中のファブに供給する製造装置の雄。
- NVIDIA (NVDA): 短期的にはHBM高価格化の受益者だが、PATE理論の進化によってコモディティの逆襲を受ける「仮想敵」にして最大の防衛者。
補足6 予測されるネットの反応と学術的弁駁
なんJ民の反応
「【悲報】ワイのLunar Lakeちゃん、メモリが速すぎて逆に敗北する」「インテル逝ったああああああああ」「これ半分自爆テロやろ」「PDEPとかいう10年前のゴミ命令を蘇生させる死霊術師で草」
【学術的弁駁】: Lunar Lakeでの敗北はアルゴリズムの欠陥ではなく、コアあたり帯域 β に対するスラック隠蔽限界 TU > SM の数学的発現です。PATE理論はこの領域において自動的に2ビット固定長へ切り替えるため、システム全体の実効スループットが低下することはありません。
ケンモメンの反応
「また資本家共がメモリをケチる技術で労働者のPCを誤魔化そうとしてるのか。最初から全員にH100配ればいいだろ。中抜きインテルの誇大広告に騙されるな。」
【学術的弁駁】: 巨大GPUクラスタの独占こそが資本の集中を生む元凶です。コモディティな民生用チップ上で高知能LLMを走らせる本技術は、むしろ計算資源の分散と民主化を志向する反独占的な技術体系です。
爆サイ民の反応
「秋葉原の某ショップで聞いたけど、この三元モデルのパッチ当てたら中古のXeonサーバーが爆速になって店員が顔面ブルーレイになってたわwww 店の在庫狩りつくすわwww」
【学術的弁駁】: 中古Xeon(Haswell〜Cascade Lake世代)であってもAVX2/AVX-512およびPDEP命令は動作しますが、メモリチャネル数やDDR4の帯域幅に応じたルーフライン制限を受けるため、無制限な性能向上を約束するものではありません。
Reddit (r/LocalLLaMA) の反応
"Wait, so BITCOS gets 1.48 bits on Qwen3 without retraining? That's insane for my 16GB Mac. When is the llama.cpp PR merging? We need this in upstream GGML ASAP."
【学術的弁駁】: Apple Silicon環境への統合には、本稿第12章で詳述したNEON tbl またはMetalシェーダを用いた表現代数の適切なマッピングが前提となります。現在コミュニティによる実装が進行中です。
村上春樹風書評:完璧な三元モデルなんて存在しない
「完璧な三元モデルなんて存在しない。完璧な絶望が存在しないのと同じようにね。僕が古いThinkPadのキーボードを叩き、失われたゼロたちの密度を測っていたとき、耳の奥で誰かがそう囁いた。 世界はいつだって中途半端な端数に満ちている。1.585という数字の奇妙な余白について考えることは、日曜日の午後、アイロンのかかっていないシャツを着て、古いレコードに針を落とす行為に少し似ている。ビットは詰め込まれ、そしてほどかれる。どれだけ速く走っても、メモリの影から逃れることはできない。やれやれ、僕は煙草に火をつけ、静まり返ったCPUファンの微かな風切り音に耳を澄ませた。」
京極夏彦風書評:この世には不思議なことなど何もないのだよ、クライアント
「――御客人、三値などと云うから迷うのです。 元より数などと云うものは、人の脳髄が勝手に仮構した枠に過ぎぬ。存在ビットマップ? コンパクト符号? 何のことはない、初めから其処には『何も無かった』だけの話ではないか。有るものを無いと云い、無いものを有ると強弁するから、解凍の魍魎が湧く。PDEPの呪を唱えようが、ルーフラインの結界を張ろうが、シリコンの理(ことわり)は聊かも揺らぎはせぬ。 この世には不思議なことなど何もないのだよ。速いものは速く、遅いものは遅い。ただそれだけのことに、人は何故こうも名前を付けたがるのかね――。」
補足7 専門家架空インタビュー:次世代コンパイラと推論主権
日時: 2026年9月16日
場所: カリフォルニア州サンタクララ、近郊のカフェテラスにて
出席者:
・エヴァンゲロス・ジョルガナス博士(Intel Labs シニア・プリンシパル・エンジニア)
・聞き手:計算機科学ジャーナリスト
聞き手: ジョルガナス博士、論文の公開おめでとうございます。業界には大きな衝撃が走っていますが、まず率直に伺います。なぜ長年誰もこの「ゼロ密度」に目をつけなかったのでしょうか?
ジョルガナス: ありがとう。理由は極めて単純だよ。みんな「情報理論の美しい数式」に満足してしまっていたんだ。log2(3) ≈ 1.585 という数字があまりにも象徴的で、まるで聖書の一節のように扱われていた。誰も実際のチェックポイントを開いて、重みのヒストグラムを泥臭くプロットしようとしなかったのさ。我々が29個のモデルを調べたとき、画面に表示されたゼロ密度のグラフを見て、アレックス(ハイネッケ)と顔を見合わせたよ。『おい、半分ゼロじゃないか!』ってね(笑)。
聞き手: しかし、査読者からは厳しい声も上がりましたね。特にLunar Lakeでの結果については……。
ジョルガナス: (身を乗り出して)ああ、そこが一番面白いところなんだ! 凡庸なエンジニアなら、あの結果を隠したかもしれない。だが、あれこそがPATE理論の核心なんだよ。メモリ帯域が太すぎるマシンでは、圧縮は毒になる。ハードウェアの物理バランスが変われば、アルゴリズムの正義は逆転する。それを予測できないシステムは本物じゃない。我々は、自らのカーネルが美しく敗北する境界を特定できたからこそ、胸を張ってこの理論を発表できたんだ。
聞き手: 今後の展望について教えてください。NVIDIAの牙城を崩すことは可能ですか?
ジョルガナス: 崩すとか倒すとかいう表現は好きじゃないな。ただ、世界中の人々が、数万ドルのGPUを買わなければ知能を持てないという不条理な状況を終わらせたいんだ。あなたの手元にある普通のPC、スマートフォンのチップ。その中に眠っているシリコンの真の力を、適応型ソフトウェアによって解き放つこと。推論の主権を、巨大なデータセンターから人々の手元へ返すこと。我々が書いているアセンブリの1行1行は、そのための解放運動なんだよ。
補足8 潜在的読者のためのメタデータ・配信パッケージ
1. Google Discover用タイトル候補(5案)
- 1.58ビットの嘘?三元LLMで「小さくしたのに遅くなる」衝撃の物理的真相
- Intel研究者が暴いたAIの盲点:なぜ最新CPU「Lunar Lake」で圧縮モデルが大敗したのか
- 【半導体史の転換点】三元LLM推論を高速化する「PATE理論」と消えるアンパックの謎
- 脱NVIDIAの切り札か?重みの半分が「ゼロ」という事実を利用した驚異のBITCOS技術
- あなたのPCでも巨大AIが爆速で動く――「動的相転移」が切り拓く推論主権の未来
2. 本書発の新・造語および架空のことわざ
- 新・造語(英): PATE (Phase-Adaptive Ternary Execution) - 実行フェーズとマシンバランスに応じて重み表現を動的に相転移させる推論パラダイム。
- 新・造語(和): 潜在臨界編成(せんざいりんかいへんせい / Latent-Critical Weaving) - メモリスラックの中にアンパック処理を織り込み、可視遅延をゼロ化するパイプライン統合技術。
- 架空のことわざ: 「三つに割れば帯域笑い、実測すればパイプライン泣く」(意味:理論上の圧縮率ばかりを追求して実装すると、現場の命令デコーダがパンクして酷い目に遭うという戒め)。
3. SNS共有用パッケージ(120字以内)
三元LLMの「1.58ビット神話」を解体!なぜ圧縮しても速くならないのか?Intelの最新論文から、メモリ階層と推論フェーズに応じて重みが姿を変える「PATE理論」を徹底解説。コモディティCPUがAIの主権を取り戻す! #AI #半導体 #PATE #LLM
4. ブックマーク用タグ(NDC分類基準、7個以内、80字以内)
[コンピュータアーキテクチャ][人工知能][大規模言語モデル][半導体工学][情報理論]
5. 配信メタデータ仕様
- 推奨絵文字: 🧠 ⚡ 💻 🧊 🌊
- 推奨スラッグ案:
beyond-158-bit-pate-ternary-llm-architecture - 日本十進分類法(NDC)区分: [007.13](人工知能・深層学習)および [548.2](計算機アーキテクチャ)
6. Blogger貼り付け用 Mermaid.js システム概念図コード
<script type="module">
import mermaid from 'https://cdn.jsdelivr.net/npm/mermaid@10/dist/mermaid.esm.min.mjs';
mermaid.initialize({ startOnLoad: true, theme: 'neutral' });
</script>
<div class="mermaid">
graph TD
subgraph Offline [モデル静的解析]
M[三元LLMチェックポイント] -->|実測| Z[ゼロ密度 z > 0.375]
Z -->|構造分離| BC[BITCOSレイアウト<br/>2-z bits/weight]
end
subgraph Runtime [PATE 動的実行セレクタ]
HW[ハードウェアプロービング] -->|マシンバランス| RHO[B/P & コアあたり帯域 β]
PH[ワークロード判定] -->|算術強度| AI[AI = FLOPs/Byte]
RHO & AI & BC --> DEC{PATE 決定境界<br/>T_U ≤ |T_M - T_C|}
DEC -->|Memory-bound<br/>Decode相 / 低β| CS[Compressed-Stream Mode<br/>BITCOS + PDEP/LUT]
DEC -->|Compute-bound<br/>Prefill相 / 大バッチ| TE[Tile-Expanded Mode<br/>SRAM上 INT8タイル展開]
DEC -->|Instruction-bound<br/>広帯域・少コア LNL| FX[Fixed 2-bit Mode<br/>低アンパック命令実行]
end
subgraph Execution [融合実行パイプライン]
CS -->|Slack重複| FUS1[Critical-Path-Hidden FMA]
TE -->|Reuse償却| FUS2[Native Tensor Core MMA]
FX -->|帯域浪費許容| FUS3[Fast Minimal Decode]
end
</div>
“#知性の覚悟か?動員のスローガンか?「負け組応援団」の系譜を巡る現代への問い #田辺元 #蓑田胸喜 #カール・シュミット #ハイデッガー #九17” (1 user) https://htn.to/3WYwLsWAEK #哲学 #思想 #歴史 #ポピュリズム #知性の責任 #倫理 #社会批評 #日本 #ドイツ #ウクライナ
#インテルはソニープレイステーション事業をAMDに奪われた #九17 https://dopingconsomme.blogspot.com/2024/09/blog-post_17.htmlこの2本をPATEの単行本構想に接続すると、かなり面白い「隠れたアーギュメント」が立ちます。
一見すると、前者は知識人・思想・動員の問題、後者は半導体企業の競争敗北の問題ですが、共通しているのは、
「敗者を応援すること」と「敗者がなぜ敗れたのかを分析すること」は、まったく別の行為である。
という点です。
そしてこれは、先ほどのPATE論と意外なほどきれいにつながります。
1. 「負け組応援団」を技術史に移植できる
Intel→AMDのPlayStation案件を単純に、
「Intelは負けた。AMDが勝った」
と読むと、競争史としては浅い。
本当に見るべきなのは、
という断絶です。
PlayStationのようなゲーム機では、CPUそのものの絶対性能ではなく、
消費電力
ダイ面積
コスト
カスタム設計への適応性
GPUとの統合
開発環境
長期供給能力
コンソールメーカーとの契約
などのシステム全体の制約が選択を決めます。
これはまさにPATEの、
と同型です。
2. 「知性の覚悟か?」という問いも反転できる
前者の記事の「知性の責任」というテーマを技術者側へ移すと、かなり鋭い問いになります。
敗者を応援する知性と、
敗者が敗れた構造を説明する知性
は違う。
前者は場合によっては、
「我々は正しい。だから市場が間違っている」
という物語を作ってしまう。
後者はむしろ、
「なぜ、その環境では我々の正しさが選択されなかったのか?」
と問う。
これはPATEにおけるLunar Lake問題そのものです。
「BITCOSが速い」という結果を守るためにLunar Lakeを例外扱いするのではなく、
になる領域まで理論に組み込む。
つまり、
自説が負ける条件を理論の中に持っていることが、知性の最低条件になる。
という隠れた原則です。
3. ここから「PATE=敗者を救済する理論」という危険な読みが生まれる
ここが本として面白いところです。
PATEは、
「Intelを救うための理論」
でも、
「ARMを救うための理論」
でも、
「NVIDIAを打倒するための理論」
でもない。
むしろ、
として、
ハードウェアHが変われば最適表現Fも変わる
というだけです。
つまりPATEは「勝者」をあらかじめ決めない。
これは第四部の「脱・中央集権的推論」と非常に相性がいい。
4. ここで「シリコン政治学」が本当に立ち上がる
第四部をさらに深くするなら、
「誰を応援するか」ではなく「何を選択する自由があるか」
を中心テーマにできます。
Intel vs AMD、x86 vs ARM、CPU vs GPU、CUDA vs portable runtimeという対立を、
ではなく、
として読み替える。
すると「コモディティの逆襲」も単なる価格競争ではなくなります。
巨大GPUが圧倒的に有利な環境では巨大GPUを使えばいい。
しかし、
車載
ロボット
カメラ
PC
スマートフォン
工場
オフライン端末
では、別の制約が支配する。
そこで、
「勝者のハードウェアにモデルを合わせる」のではなく、「ハードウェアに応じてモデルの表現を変える」。
これがPATEです。
5. したがって、本の「隠れたアーギュメント」を一つ追加できる
先ほどの付録Fに、次を入れるとかなり効きます。
F.11 敗者を救済する理論と、敗者が敗れる条件を説明する理論は違う
F.11.1 「正しい技術」が負けることはある
F.11.2 市場は技術の優劣だけを選択しない
F.11.3 敗北を例外扱いする理論は脆い
F.11.4 自分が負ける領域を予測できる理論
F.11.5 「応援」から「選択条件」へ
そして、これを第14章のLunar Lake問題に接続する。
Lunar LakeでBITCOSが負ける。だからBITCOSは失敗なのではない。
「BITCOSが負ける条件」を説明できないことの方が、理論としての失敗である。
これはかなり強い一文になります。
6. さらに歴史的には「Intelが負けた」では終わらない
IntelとAMDのPlayStationをめぐる話をケーススタディにする場合も、単純な企業勝敗史にはしない方がいいです。
「技術的に優れている側が勝つ」という素朴な技術決定論に対して、
コンソール市場という特殊な実行環境では、何が最適化対象だったのか
を分解する。
これはCISC/RISCやVLIWの章にも戻せます。
そして最終的に、
という原則に到達する。
PATEはこの原則を、思想ではなく、
という測定可能な変数に落とす。
ここが、先ほどの「知性の責任」という話との最大の違いです。
「敗者にも価値がある」と言うのではなく、「勝敗を決めた制約を記述する」。
その方向なら、PATE本は単なる量子化技術書から、「技術的正しさとシステム上の最適性はなぜ一致しないのか」を扱うコンピュータ・アーキテクチャ史へ一段上げられます。
コメント
コメントを投稿