OpenAI の新チップ「Jalapeño」が Nvidia GPU の欠点を克服する設計へ

2026年8月30日 08:51 浙江

記事：渣B

TL;DR

もし私が大規模モデルチームのインフラ責任者として生まれ変わったら、Nvidia の GPU やシステム全体に対して不満を抱き、自社のチップ開発に着手するだろう。

Nvidia のネットワーク部分については論じるまでもない。RoCE にはさまざまな欠陥があり、DPU も性能やセキュリティ面で多くの課題を抱えている。そのため今年 Nvidia が「Scale-in」を強調し始めたのも無理はないが、クラウド経験の乏しいチームがこの分野で成熟するには最低でも 5 年はかかるだろう。

具体的な例として、最近 MaaS のオンライン業務における PD（Prompt Dispatch）分離シナリオで、転送プロトコルを CIPU の eRDMA に置き換えたところ、TTFT（Time to First Token）が 40% 短縮された。これは TTFT の SLA を維持したまま Prefill 計算サーバーを大幅に削減できることを意味する。

GPU マイクロアーキテクチャについては、昨年『Inside Nvidia GPU: 谈谈 Blackwell の不足并预测一下 Rubin の微架构』（英語版：https://github.com/zartbot/blog/issues/3）という記事で Blackwell の欠陥を詳しく批判した。この内容は黄仁勲氏（ジェンソン・フアン）さえも Nvidia 内部で共有していたほどだ。

当時は SM（Streaming Multiprocessor）にスケジューリング用のスーパースカラコアを追加するよう提案したが、それが OpenAI の Jalapeño で実現された。もう一つの課題として、私は長年 Nvidia に指摘してきたが、GPU アーキテクチャでは L2 キャッシュがチップ面積の多くを占有しているにもかかわらず、制御が難しく、キャッシュ階層構造も厚すぎる。UMA（統一メモリアーキテクチャ）を維持するために L2 の 2 つのパーティション間で 200〜400 クロックサイクルの遅延が発生し、外部 I/O との相互作用も直接 SMEM（Shared Memory）に書き込めないという問題がある。OpenAI はこの課題に対し、Jalapeño で L2 キャッシュを完全に削除することで解決した。

また数ヶ月前、Anthropic の面接試験で出題された問題も、実際には自社開発のトレーニング・推論チップ関連の業務準備の一環であり、社内ネットワーク上でさまざまなアクセラレーターアーキテクチャとその設計手法について詳細に分析していた。

今回の Jalapeño 分析記事では、私が大規模モデルチームのチップ責任者になったと仮定し、モデル要件と現在の Nvidia GPU の欠陥を起点に、第一原理からどのように設計すべきかを段階的に導出する。同時に、社内での設計プロセスも再現していく。この手法を通じて、Jalapeño のチップアーキテクチャとその背後にあるツールチェーンについて解説する。

本記事の構成は以下の通りだ。

1. Nvidia GPU から考える

1.1 推論（インフェレンス）における理論上のピーク性能

1.2 アーキテクチャが E2E レイテンシを決定する

1.3 Nvidia GPU の欠点

1. 推論用チップの設計方法

2.1 推論用チップの第一原理

2.2 推論インフラの現状

2.2.1 各段階における性能要件

2.2.2 KVCache がもたらすコスト

2.3 推論用チップのアーキテクチャ

2.3.1 プログラミングインターフェースの変化

2.3.2 アーキテクチャ設計におけるトレードオフ

1. OpenAI の「Jalapeño」チップアーキテクチャ

3.1 チップアーキテクチャ概観

3.1.1 Core Slice アーキテクチャ

3.1.2 NOC（オンチップネットワーク）アーキテクチャ

3.2 Core Slice アーキテクチャの詳細解説

3.3 メモリサブシステム

3.4 オンチップネットワーク

3.6 システムアーキテクチャ

1. ソフトウェアアーキテクチャ

1. チップ設計プロセス

5.1 重要な前提：言語を先に刷新し、その上で AI を投入する

5.2 デザインのクローズドループ形成

1. 将来展望

6.1 Nvidia GPGPU アーキテクタとの比較

6.2 OpenAI のロードマップ

6.3 分析と総括

⚠️ より詳細な評価や、設計プロセス全体を再現する試みについては公開できません。微信公众号（WeChat Official Account）のレイアウトは PC での閲覧に適さないため、Github に再構成したバージョンを掲載しました。

中文版：https://zartbot.github.io/blog/arch/jalapeno/index.html

英文版：https://zartbot.github.io/blog/arch/jalapeno/en.html

1. Nvidia GPU から考える

まず思い浮かぶのは、現在の状況を象徴するあるジョークです。

製品管理担当者がこう言います。「隣の NV（Nvidia）がやっているのだから、なぜ私たちはやらないのか？」

また別の製品管理担当者はこう反論します。「隣の NV すらやっていないのに、なぜ私たちができるというのか？」

さらに別の製品管理担当者はこう考えます。「隣の NV がやっているなら、なぜ私たちもやる必要があるのか？」

そして最後に、「隣の NV がやっているのだから、なぜ私たちは手を出せるのか？」と不安がるのです。

加速器メーカーの多くは、Nvidia のやり方を模倣することに夢中です。それを「CUDA エコシステムの互換性」と称して。しかし、どんなに真似をしても、全体として Nvidia よりも 2〜3 世代遅れになるのが実情です。

一方、極端なアプローチを取る企業もあります。奇抜なアーキテクチャを採用して、無理やり追い越そうとするのです。しかし、ほとんどのチップメーカーが Nvidia と決定的に違うのは「フルスタック能力」の欠如です。モデル構造やアルゴリズムから始まり、ソフトウェアエコシステム、そして最終的なチップ実装に至るまで、どこでトレードオフ（妥協）が可能なのか、そのトレードオフによって何を勝ち、何を失うのかを明確に説明できる人はほとんどいません。

私自身は、アルゴリズムからチップ設計に至るまでのフルスタックを理解しているため、この機会にチップ設計の全体像と Nvidia GPU の欠点を包括的に解説します。

1.1 推論（インフェレンス）における理論上のピーク性能

GPU は長年、「多数の並行する Warp を活用してメモリアクセスの遅延を隠蔽する」という仕組みで大量データを処理してきました。しかし、この仕組みは本質的に大規模なバッチサイズに依存しています。

バッチサイズが小さくなると、スケジューリング可能な Warp の数が減り、パイプラインが十分に埋まりません。その結果、性能の理論的上限（Roofline）から遠ざかってしまいます。これが推論シーンで頻繁に発生する問題であり、遅延隠蔽メカニズムが機能しなくなっている状態です。

推論フェーズは本質的にメモリバンド幅制約（Memory-bound）であるため、OpenAI は非常にシンプルな計算を行いました。

仮にスケールアップドメイン内に N 個のチップがあり、各チップの HBM バンド幅が M であるとします。また、モデルのパラメータ規模を K としましょう。計算自体にはボトルネックがないと仮定すれば、理論上、出力されるトークンの最大速度は次のように見積もることができます。

image 理論上、推論（デコーディング）の最適化を行わなければ、1 トリオンパラメータを持つモデルはユーザーあたり秒間 1,000〜2,000 トークンの処理速度を達成できるはずです。しかし現在、B300 などのプラットフォーム上で動作する大半の推論フレームワークでは、最大でもユーザーあたり秒間 100〜200 トークン程度しか出ず、理論性能のわずか 10% に留まっています。

つまりこのページで言いたいのは、HBM（ハイバンド幅メモリ）の帯域を理論限界まで引き上げても、レイテンシにはある程度の下限が存在するということです。単に帯域を増やせば低遅延が解決するというわけではないのです。帯域だけが変数ではない以上、問題はアーキテクチャそのものにあるはずです。

1.2 アーキテクチャが E2E レイテンシを決定する

約一年前の記事で私は、プロセッサのワークロードという観点から、現代のコンピュータアーキテクチャにおける CPU と GPU の設計哲学は根本的に異なることを示唆しました。GPU は膨大な数の並列スレッドを用いて待ち時間を隠蔽し、スループットの最大化を図る一方で、CPU はチップ面積やトランジスタ消費電力といったコストを惜しまず、微細なアーキテクチャの工夫を尽くして単一スレッドの実行遅延（Single-Thread Latency）を極限まで圧縮することに注力しています。

image CPU は分岐予測によって処理の先行きを探り、マルチディスパッチとアウト・オブ・オーダー実行で隙間を縫って並列計算を行い、多段キャッシュによってデータへのアクセス距離を短縮します。さらにハードウェアレベルでのキャッシュ一貫性を活用し、コア間のナノ秒単位のローカルなデータ共有やメモリ同期の不要化を実現することで、単核計算、メインメモリアクセス、そしてマルチコア間での協調処理に伴う待ち時間を巧妙に隠蔽しています。

一方、GPU はトランジスタ予算のほとんどを密な演算論理ユニット（ALU）へ割り当て、SIMT 構造を用いて数万規模の並列スレッドを組織します。ハードウェアレベルで超高速なスレッドコンテキスト切り替えを行い、あるスレッドがメモリアクセスを待っている間に即座に他の準備完了スレッドをスケジューリングする仕組みを持っています。しかし、GPU は単一命令の処理速度には一切関心がなく、圧倒的な並列度によってデータフロー全体の処理遅延を、膨大な計算タスクの中に埋没させてしまいます。

LLM（大規模言語モデル）の推論においては、この状況は以下の表の A* の位置に該当します。バッチサイズが 1 の場合、1 レイヤーで数回の GEMV（行列積演算）が発生し、さらにミリ秒単位で数十層を通過する必要があります。また、各レイヤー内でも大量の並列計算後の結果同期を行わなければなりません。

• 遅延敏感：遅延不敏感
• データ規模大：A* (LLM) / GPU
• データ規模小：CPU / NA

同時に、OpenAI が HotChips のセッションで示した資料にも、この問題について言及されています。私たちが実際に測定するレイテンシの多くは、実は基盤となるアーキテクチャそのものによって決定されているのです。

image図に示される「長い遅延経路→演算子の到着遅れ→計算ユニットのブロック」は、従来の Roofline Model に新たな次元を加えるものです。GPU チップにおいて、ワークロードを単に「計算バウンド（Compute bound）」と「メモリバウンド（memory bound）」で分類するだけでは不十分です。「ブロック」という要素も考慮する必要があります。この状態では、リソースのいずれも最大利用率に達していませんが、すべてのリソースが待機している状況です。結果として利用率は低下しますが、性能を制限する上限（天井）はまだ遥か彼方にあります。

1.3 Nvidia GPU の欠陥

ここでは、OpenAI が提示した見解について詳しく掘り下げていきます。

まず、OpenAI が指摘する「統一メモリサブシステムでは経路が過度に競合しやすい」という点から議論を始めましょう。Nvidia のアーキテクチャは、統一された L2 キャッシュとグローバルアドレス空間を採用し、どのコアも任意のアドレスに等価にアクセスできます。この統一アクセス方式はプログラミングの複雑さを隠す一方で、多くのエンジニアリング上の制約を生み出しました。

例えば、SM（Streaming Multiprocessor）の数が増えるほど帯域幅への要求が膨大になります。Ampere アーキテクチャ以降、L2 は 2 つのパーティションに分割されました。Hopper ではパーティション間をまたぐアクセスで 200 サイクルもの遅延が発生します。Blackwell や Rubin ではさらに深刻な影響が出ます。Blackwell の各 Die には 2 つの L2 パーティションがあり、Dual-Die 構成では合計 4 つになります。Blackwell ではこの遅延が約 400 サイクルに達し、Rubin では HBM4 の累積帯域幅がすでに L2 バンド幅に近づいているため、NV HBI（Nvidia High Bandwidth Interconnect）の影響も受けます。私の見解では、統一メモリアクセス方式により、HBM の 22TB/s という理論上の実効帯域幅利用率は実際には 60%〜70% に留まると考えられます。さらに、並行リクエストがポートの処理能力を超えた場合、発生するキューイング遅延（Queuing Delay）のコストは、データ転送そのもののコストを遥かに上回ります。

もう一つの課題は、データ同期やワープ（warp）スケジューリングにおける問題です。現代の GPU アーキテクチャでは、各計算コアが独立したプログラムカウンター（PC）を持ち、非同期で自由に進行します。また、TMA（Tensor Memory Accelerator）などの非同期メモリアクセスデバイスによってスループットは向上しました。しかし、計算コアが独立しているがゆえに、ハードウェア側が「誰がいつ到着するか」を予測することができません。その結果、グローバルフェンスを実行する際には、実際にすべてのコアに「問いかけ」を行う必要があり、場合によってはシステムレベルの ScaleUP/ScaleOut ネットワーク全体でバリア（barrier）処理を行う必要があります。これが OpenAI が指摘した、「独立したコアが同期を外れると、グローバルメモリフェンスのコストが高くなる」という現象です。

最後に、ネットワークアクセスを仲介する集中リソースについて言及します。数百から数千のコアが外部ネットワークへ通信を发起する際、集中型の DMA コントローラーや MMIO インターフェース、あるいは共有の NIC ポートを経由してデータを送受信すると、仲裁ロジック（Arbiter）が深刻なスループットのボトルネックとなり、レイテンシのホットスポット化します。その結果、ハードウェア本来の高速転送能力が完全に阻害されてしまいます。

これは NOC、ScaleUP、ScaleOut といった複雑な相互接続システムにまつわる一連の問題です。具体的には、KVCache のコピー時に用いられる CopyEngine や、ScaleOut NIC における doorbell の制約などが挙げられます。また、多くの通知を直接 SMEM（Shared Memory）へ書き込めないため、Remote GMEM への書き込みを余儀なくされ、そこから CUDA Core が GMEM をポーリングするという非効率なフローが発生します。

こうした背景から、私は過去数年にわたり、Nvidia に対して「SM 内にスーパースケーラ型コアを設け、プログラマーが命令の発行とスケジューリングをより細かく制御できるようにすべきだ」と提言してきました。同時に、L2 キャッシュにおいてもユーザーにより多くの制御権限を与え、SM 内の SMEM とさらに密接に融合させるべきだと主張しています。ちょうど OpenAI の Jalapeño も同様の選択を行いました。次章では、推論業務の第一原理からアプローチし、推論に適したチップをどのように設計すべきかを順を追って解説していきます。

1. 推論用チップの設計方法

前節で指摘した欠陥を仔細に分析すると、それらはすべて「汎用性」に対する代償であることがわかります。ここで強調しておきたいのは、GPGPU の設計が劣っていると言っているわけではありません。当時の各決定は正しく、それが Nvidia の成功の鍵でした。CUDA SIMT というプログラミングモデルは、過去 10 年以上にわたり並列計算プログラムの記述を劇的に簡素化してきました。

しかし、LLM（大規模言語モデル）の推論という特定のシナリオにおいては、この前提が数多くの問題を引き起こしています。さらに Coding Agent の発展に伴い、Auto-Kernel（例：KDA など）が百花斉放しており、「汎用性に対する代償をこれほど高く支払う必要はないのではないか」という声も聞こえてきます。

しかし、これはあくまでチップのアーキテクチャという視点に過ぎません。より重要なのは、推論ビジネス全体を「第一原理」として捉え直すことです。

2.1 推論用チップの第一原理

第一原理に基づけば、推論サービスの経営視点における主目的は、インフラコストとサプライチェーンの制約の中で、ユーザー体験の SLO（サービスレベル目標）を満たしつつ、可能な限り多くのトークンを出力することです。特に北米におけるサプライチェーン最大の制約要因は、電力供給の限界である可能性が高いです。

そのため、OpenAI の講演資料では、「使用チップ数」「チップあたりのスループット」「TTFT（Time to First Token）」を主要な目標指標として掲げるのを意図的に避けています。代わりに選定されたのは「リクエストレベルでのユーザー体験」と「リクエストあたりのエネルギー消費量」です。

image 并发処理によって一部の待ち時間を隠蔽することは可能ですが、単一タスク内の依存関係チェーンを無理に並列化することはできません。そのため、全体のリクエスト完了時間を示す「TTLT（Time to Last Token）」と、ストリーミング的な対話速度の近似値である「TBT（Time Between Tokens）」という指標を用います。

しかし、バッチサイズを小さくすると TBT は改善されるものの、重みやネットワークの再利用効率が低下し、トークンあたりのエネルギー消費が増加することが分かっています。逆にバッチサイズを大きくすれば tokens/s/kW の数値は向上しますが、リクエストの待ち行列が長くなり、ユーザーごとのトークン間隔も延びてしまいます。つまり、「応答遅延（TTLT）」と「エネルギー効率（tokens/Joule）」の間にはトレードオフの関係があり、パレート最適解が存在します。

これを形式的に記述すると以下のようになります。

LLM の推論は、実質的には CPU 業界が過去 30 年かけて取り組んできた「単一スレッドの遅延問題」そのものです。

CPU 界の知見と LLM 推論への対応

• 分岐予測（Branch Prediction）：次に来る命令を予測し、間違えばロールバックする → 投機的デコーディング（Speculative Decoding）：次の数トークンを予測し、間違えば破棄する。
• 乱序実行（Out-of-order Execution）：依存関係チェーンの隙間を見つけて独立した処理を実行する → 乱序コア。
• キャッシュとプリフェッチ：遅延をクリティカルパスから外す → ハードウェア管理 L1 キャッシュと依存関係を考慮したプリフェッチ。スクラッチパッドや DMA に頼るのではなく、これらを活用する。
• SMT（スレッドレベル並行処理）：別のスレッドで隙間を埋める → 別のシーケンス（リクエストレベルのバッチ処理）で隙間を埋める。

加えて、「命令レベル並行性（ILP）の上限」により、実行幅を広げても効果は頭打ちになります。同様に、「投機的深度」を増やしても「受容率の上限」に達すると、それ以上の効果は見込めません。

このアーキテクチャの視点から導き出される結論は非常にシンプルです。

1. データ・ローカリティ（Data Locality）を最大化し、データ転送のコストを最小化する。
2. エンドツーエンド（E2E）遅延の観点から、命令実行の並列性をさらに高め、アイドル状態での待ち時間を減らす。

次のセクションでは、推論技術の詳細について掘り下げて分析していきます。

2.2 推論インフラの現状

#### 2.2.1 各フェーズにおける性能要件

一般的な推論システムでは、主要な業務処理は「Prefill（初期化）」「Decode（生成）」「SpecDecode（投機的デコーディング）」の 3 つに大別されます。従来の定性的な見解では、「Prefill は計算集約型（Compute Bound）、Decode はメモリ集約型（Memory Bound）」とされてきました。

しかし、E2E 遅延という視点で考えると、SpecDecode のフェーズには「Draft モデルによる推測」と「Verify（検証）」の 2 つの段階が含まれており、これらはそれぞれ Roofline モデル上の異なる位置に存在します。

OpenAI は HotChips のセッションでもこの問題に触れましたが、その説明は不十分でした。例えば、3 つ目のスライドではメインタイトルが「Spec-verify」、サブタイトルが「Decode」となっていますが、自己回帰的な Decode とバッチ処理による Verify は Roofline 上で全く異なる位置にあり、演算強度（Arithmetic Intensity）の差は数桁にも及びます。

imageこの図は、分析を簡略化するために 1 つのリクエスト処理に焦点を当てた場合の、推論チップの 4 つの主要なフェーズを示しています。

1. Prefill（事前填充）: この段階では、入力コンテキスト全体が一度に処理されます。隠れ次元が D で、シーケンス長が L の一般的な Transformer モデルにおいて、投影層と FFN（フィードフォワードネットワーク）の主要な計算量は D に比例して増加し、フルアテンションのスコア部分は L^2 に比例します。L が十分に大きくなると M 次元が確保され、重みが多数のトークン間で再利用されるため、Tensor Engine は高い利用率を達成しやすい状態になります。つまり、これは「Compute-bound（計算集約型）」な処理であり、計算量は高く、メモリ帯域幅は相対的に低いことを意味します。通信面では、Collective 操作のメッセージサイズが L に比例して増加するため、遅延ではなく帯域幅がボトルネックとなり、スムーズなスケジューリングが可能になります。

1. Decode（デコード）: 自己回帰モデルはこの段階で、各シーケンスに対して 1 トークンずつ逐次処理を行います。つまり、L=1 の状態です。行列計算は GEMV（General Matrix-Vector Multiplication）に退化し、結果として 1 つの重みベクトルが 1 行分のデータのみを処理することになります。短〜中程度のコンテキストでは通常、Weight HBM（High Bandwidth Memory）の帯域幅がボトルネックとなりますが、長文コンテキストになると次第に KVCache の HBM 帯域幅への依存が高まります。分散システムにおいては、各レイヤー境界で依然として Collective 操作が発生しますが、そのメッセージサイズは極めて小さいため、ネットワーク同期の遅延が最大のボトルネックとなる可能性があります。

1. Draft Model（ドラフトモデル）: このフェーズはモデルの種類に大きく依存します。従来の自己回帰型ドラフトモデルや EAGLE 系手法では、複数の逐次的なドラフトステップを実行する必要があり、重み/KVCache の帯域幅と同期遅延に制約されることが多いです。一方、DFlash はブロック並列拡散バックボーンネットワークの単一フォワードパスでドラフトブロック全体を生成します。DSpark は並列バックボーンの前向伝播を 1 回実行した上で、低コストなマルコフまたは RNN シリアルヘッドを追加する方式です。DFlash や DSpark においては、ドラフト重みはバックボーン呼び出し時に読み出されるだけで済み、計算強度（Arithmetic Intensity）が L の増加とともに向上します。これによりボトルネックは行列演算の算力へとシフトし、隠れ特徴量や KV の注入帯域幅、Vocabulary Projection、あるいは軽量なシリアルサンプリングの遅延が新たな課題となります。

1. Verify（検証）: 物理的なトークン検証数は、必ずしも固定値 L に等しいわけではありません。線形ドラフトでは通常 1 つの前缀（プレフィックス）のみを検証しますが、木構造ドラフトではすべてのノードを検証する必要があります。DSpark では、信頼度とリアルタイムのハードウェア容量に基づき、各リクエストに対して検証長を動的に選択します。稠密モデルであれば、重みは今回の物理検証トークン間で再利用できますが、MoE（Mixture of Experts）目標モデルの場合、トークン予算に応じて重複排除された後の専門家（Expert）数が増加するため、重みの再利用性が大きく低下する可能性があります。効率的な検証アテンションを実現するには、1 つの前缀 KV ブロックを、同一リクエスト内のすべての候補クエリまたは木構造ノードで共有・再利用する必要があります。

これらの特徴をまとめると以下の表のようになります：

チップの計算、メモリ、ネットワーク帯域幅、同期遅延に対する要件は、処理段階によって大きく異なります。各段階に専用チップを割り当てると、リソース配分に問題が生じる可能性があります。OpenAI もこの課題について言及しています。

OpenAI は、現在の PD（Prefill-Decode）分離や AFD/Speculative Decode といった LPU（Language Processing Unit）の導入に伴う課題を指摘しています。元の資料では独自の三分法を採用しているため、prefill、draft、verify の 3 つの比率のみを示されていますが、実際には前述した 4 つの段階をそれぞれ専用のハードウェアに割り当てることで、各チップが自らのカーネルに対して最適化を図れるように見えます。

しかし、問題はそのような構成では、導入後のデバイス台数が離散的かつ相対的に固定されてしまう点にあります。4 つのプール（prefill、decode、speculate、verify）それぞれの処理能力をそれぞれとし、到着する作業量をそれぞれとします（添字は順に prefll、decode、speculate、verify に対応）。この場合、システム全体のサービス可能速度は最も逼迫しているプールの容量によって制限されます。

例えば、Speculative Decode の受容率が動的に変化することの影響を検討してみましょう。これらは互いに増減し、そのバランスを動かすのはリクエストの性質ではなく、出力トークンの経路選択です。1 つのリクエストから生成される出力トークンは、純粋な decode パスを経るか、Draft + Verify のパスを経るかのどちらかしかありません。ここで、出力トークンの中で Speculative Decode 経路を通って提出された割合をとし、1 秒間に実際に提出された出力トークンの総数をとします。

このとき、以下の関係が成り立ちます：

ここで、は 1 ラウンドの Draft モデル実行にかかるコストです。自回帰生成ではステップ数に比例しますが、ブロック並列処理では 1 回の広幅パスとして扱われるため、は物理的なトークン予算に比例します。この値はリクエストの混合比率には依存しません。

なぜなら、受容率が低下すると（＝）は 0 に近づき、高負荷下でスケジューラがプレフィックスの生存確率に基づいて低信頼度のサフィックスを削除しても同様に 0 に押し下げられるからです。逆に、受容率が回復すれば再び 1 に戻ります。

したがって、同一のリクエスト群・同一モデル・同一コンテキスト長という条件下では、decode プールと verify プールの負荷は互いに逆の極端な方向へ向かう可能性があります。decode プールは最大値に備えてリソースを確保する必要があり、draft と verify のプールも同様に最大値に備える必要がありますが、これらのピークが同時に到来することは決してありません。つまり、リソースの割り当ては各ピークの合計に基づいて行われながら、実際にはその中のいずれか一方の負荷パターンしか利用されないことになります。

異なるリソース軸に 4 つの制限項目が配置されます。prefill（初期処理）は行列演算能力を要し、decode（推論生成）は HBM の重み帯域と層ごとの同期遅延に依存します。自己回帰によるドラフト作成ではドラフト用の重み帯域と二次同期が発生し、verify（検証）では重複排除された Expert への重みトラフィックと EP all-to-all のバーストが問題となります。

つまり、これら 4 つの専用チップは単にサイズが違う同一モデルではなく、計算能力・帯域幅・ネットワーク比率がそれぞれ異なる別々のチップです。受け入れ率が上昇すれば verify の回数は減りますが、コンテキスト長が伸びれば prefill と KV コストが増加します。また、Expert 占比の高い MoE モデルへ変更すると、verify におけるネットワークと HBM の負荷がさらに増大します。

あるプールがボトルネックになると、他の 2 つのプールもパッケージ、HBM、I/O、ネットワーク、冷却に関するベースラインコストを払い続けなければなりません。ここで見落としがちなのが、運用担当者が設備をアイドル状態のプールからボトルネックのプールへ移しても、実際には使用できないという事実です。prefill チップは HBM と計算能力の比率が decode 用に設計されていないためです。

固定されたデバイス台数の割合は、同時に 5 つのリソース次元（計算、帯域、ネットワークなど）の比率を固定することと同義です。

2.2.2 KVCache のコスト

従来の推論デプロイでは、PD（Prefill-Decode）分離方式を採用することが一般的ですが、これに伴う大量の KVCache 転送は膨大な電力消費とネットワーク同期の損失をもたらします。OpenAI は HotChips で以下のように述べています。

このスライドは、アーキテクチャの選択を二択に圧縮しています。一つ目は「システム間での異種化」です。これは KV が異なるステージを移動する方式で、現在の PD 分離アプローチがこれに該当します。しかし、各段階におけるネットワークと同期の影響、そして KVCache の転送に伴う電力消費は甚大です。

二つ目は「チップ内部での異なるステージ間の KVCache 移動」です。これは、アクティブなリソースの割合がステージごとに変動することを意味します。

通常、KVCache のサイズはシーケンス長に比例して増大します。Linear Attention などのアルゴリズムでさえも、複数回の相互作用において各ラウンドのスナップショットを保持する必要があります。データの転送は ScaleUP や ScaleOut、そして NOC（Network on Chip）が既存の集合通信性能に与える影響を引き起こします。さらに、Prefill ノードの HBM から KVCache を読み出し、Decode ノードの HBM に書き込むというプロセスも、既存の推論オペレータの実行に影響を及ぼします。

具体的な分析は、「時間比」「ラウンド境界」「HBM 容量」「トランザクションセマンティクス」の 4 つの観点から定量的に展開することが可能です。

この部分は機密事項に該当するため、詳細は公開できません。しかし、記事の冒頭で触れた通り、MaaS サービスの転送プロトコルを CIPU eRDMA に変更した際、推論エンジン側のコード変更は一切行われませんでした。その結果、TTFT（First Token Time）が 40% 低下しました。

最後にチップ設計についてですが、複数の独立したアクセラレータを使用する場合、KV キャッシュはデバイス間を移動させる必要があります。一方、Jalapeño の戦略は、この異種計算能力を単一のバランスの取れたチップ内部に統合することです。

• 現在のフェーズでテンソル演算が必要であれば、より多くの計算ユニットを活性化します。
• HBM（High Bandwidth Memory）がボトルネックとなる場合は、ローカルメモリ経路を重視します。
• 通信集約型の処理が必要な場合は、collective 操作と scale-up ネットワークを活用します。
• 一時的に不要なモジュールはパワーゲーティングで停止させます。
• KV キャッシュは可能な限りローカルに保持します。

2.3 推論チップのアーキテクチャ

数ヶ月前、David Patterson は『Challenges and Research Directions for Large Language Model Inference Hardware』[1]という論文を発表し、推論ハードウェアのアーキテクチャについて議論しました。メモリサブシステムに関しては、HBM や 3D-DRAM、HBF（High Bandwidth Fabric）の使用など多くの議論がありますが、ここでは問題の簡略化のために、一部の技術的詳細には立ち入りません。

例えば、新しい HBM を活用して PNM（Processing Near Memory）やカスタムベースドダイを構築し、メモリコントローラと HBM PHY を計算用ダイから分離することで、より多くの演算リソースを確保する方法があります。また、AMD が採用している XCD（Xeon Compute Die）と FCD（FPGA Compute Die）の 3D パッケージングのように、XCD の歩留まりを向上させ、L2 キャッシュやノードオンチップ（NOC）による相互接続を拡大する手法もあります。

これらの話題は一旦脇に置き、計算用ダイ内部のマイクロアーキテクチャと相互接続アーキテクチャに焦点を当てて議論を進めます。

社内の詳細なドキュメントでは、各種演算チップのマイクロアーキテクチャについて分析がなされています。対象には Nvidia の GPGPU / LPU、Google TPU、そして Huawei Ascend 950 が含まれます。また、エージェントベースのコンパイラに関する取り組みについても一連の探索が行われています。

#### 2.3.1 プログラミングインターフェースの変化

数ヶ月前、私は Anthropic 社の採用試験で出された性能エンジニアリング関連の課題に取り組みました。この課題は、https://github.com/anthropics/original_performance_takehome にコードが公開されています。

内容は、カスタム VLIW + SIMD プロセッサ上での性能最適化です。Claude-Op モデルを使用して最適化を進め、2 位との差を 2% 以内まで縮めることができました。しかし、計算リソースとトークン費用の制約からそれ以上は進められず、結果として現在も 3 位のままとなっています。

このテーマを通じて、私はあることに気づきました。これは推論と学習のための専用 ASIC チップを設計するコンパイラシステムです。従来の CUDA エコシステムや SIMT プログラミングインターフェースは人間にとって親和性が高いものの、その代償として膨大なリソースコストが発生していました。

しかし、「エージェントをコンパイラとする」アプローチや Ligeng 氏の KDA（Kernel Dispatch and Allocation）のような研究が進展する中で、カーネル最適化はもはやエージェントに任せることが可能になりました。これにより、人間の認知上の制約に合わせてハードウェアを調整する必要はなくなり、ハードウェアの性能を最大限に引き出す時代へと移行します。

その結果、従来の SIMT の抽象化や GPGPU が推論シーンで持っていた優位性は失われます。通常、モデルアーキテクチャは事前学習段階で確定しており、推論プラットフォームに対してエージェントベースの最適化を行うための 2 ヶ月という十分な時間を確保できます。これが OpenAI の「Jalapeño」アーキテクチャが解き放った新たな自由です。

2.3.2 アーキテクチャの取舍

昨年のブログでも触れた通り、Nvidia の Warp スケジューリング機構やレイテンシ隠蔽機能は、推論シーンには適していません。また、『GPU のメモリモデルとネットワーク設計について』という記事で指摘したように、Nvidia GPGPU の複雑なメモリセマンティクスには問題があります。データパス上には Genric Proxy や Async Proxy が存在し、ScaleOut を行うためには追加の同期メカニズムが必要です。さらに、NOC（Network on Chip）やメモリ階層構造も肥大化しています。

Warp スケジューリングの仕組みを見ても、私の期待に応えるものではありません。データ転送のたびに明示的なオーケストレーションが必要であり、そのオーケストレーション自体にコストがかかります。DMA 記述子の構築、起動、完了待ちのための Barrier などです。

仮にある Tile の計算が 500 サイクルで完了するとします。しかし、DMA オーケストレーションと Barrier のオーバーヘッドが 100 サイクルかかる場合、固定オーバーヘッドは全体の 20% を占めます。一方、GPU や TPU の典型的なシナリオ（大規模バッチ処理、Tile あたり数千サイクルの計算）では、同じ 100 サイクルのオーバーヘッドは全体の 2% に過ぎず、無視できるレベルです。

不幸にも、推論のデコードフェーズはこの前者のケースに該当します。バッチサイズが小さく、各 Tile の計算量も少ないため、同期オーバーヘッドは「無視できるもの」から「支配的な要因」へと変化します。この状況では、同期を「明示的な Barrier」から「ハードウェア依存追跡（アウト・オブ・オーダー実行＋キャッシュ）」へ切り替えることで、明確な利益を得ることができます。

そこで私は、SM（Streaming Multiprocessor）内にスーパースケーラコアを追加してレイテンシをさらに制御すべきだと考えていました。今回の OpenAI Jalapeño に採用された Out-of-Order Core は、まさにこの判断を検証するものとなりました。

また、この数ヶ月は Blackwell シリーズについても、命令レベルのサイクル計測とチップ全体の分析を徹底し、『Dissecting SM_120 through Microbenchmarking』[2] という論文として公開しました。その結果、B300 のオンチップネットワークや L2 キャッシュ、そして電力管理において多くの課題が浮き彫りになりました。

一言で言えば、メモリ階層構造が重すぎます。オンチップネットワークにおいては、L2 の帯域幅を維持するために、当初から 2 つのパーティションに分割される設計になっていました。Nvidia は UMA（統一メモリアーキテクチャ）のプログラミングモデルを継続し、異なる L2 パーティションへのアクセス遅延差を隠蔽しようとしていますが、Multi-Die 構成の Blackwell や Rubin ではその代償が甚大です。

Multi-Die の構造により単一の GPU に 4 つの L2 パーティションが存在するにもかかわらず、依然として UMA を維持しようとすると、推論フェーズにおけるメモリアクセスの損失は数百サイクルに達します。これは明らかに得不偿失（得るものより損うものが大きい）と言える状況です。

もう一つの課題は、RDMA などのシナリオです。ここでは多くの通知メカニズムを L2 をバイパスし、直接 SMEM（Shared Memory）へ書き込むことを望んでいます。そうすれば、SM は SMEM をポーリングすることで、より予測可能な遅延で情報を取得できます。一方、現在の GMEM のポーリングでは、Warp Scheduler が Yield してしまい、さらに大きな遅延を招くリスクがあります。

また、L2 キャッシュの動作メカニズムは我々にとって制御が困難です。重要な同期用のカウンタが、SM からの大量なメモリアクセスによって Evict（キャッシュから追い出される）されてしまう可能性も十分にあります。私はかねてより Nvidia に対して、L2 キャッシュの一部を SMEM のように自分で制御できる領域として割り当てる仕組みについて提案し続けています。

さらに深刻な問題もあります。同期メカニズムの影響により、チップ内（intra-chip）だけでなく、チップ間（inter-chip）の集合通信性能も遅延によって制約を受けます。これが原因で、多くのシナリオにおいて TP（Tensor Parallelism）並列の規模をこれ以上拡大できず、ボトルネックとなっています。

これらの問題はすべて UMA の影響に起因しているように見えます。そこで自然な発想として、「NUMA 構造を採用し、集合通信に関連する重要な処理を NOC（ネットワーク・オン・チップ）やインターコネクト（ScaleUP/Scale Out）で十分に最適化できないか」という問いが生まれます。

もちろん、Nvidia が商業企業としてピークスループットを追求するのは理解できます。しかし、推論システムという観点からは、目標関数に基づいて 2 つの独立した結論が導き出せます。それぞれが微細なアーキテクチャとシステムアーキテクチャに該当します。

推論 A：有効なトークンは、「低遅延下での高利用率」によってのみ達成可能であり、バッチ数を積み上げることで得られるものではありません。具体的には以下の通りです。

• メモリへのアクセス経路を短縮する：メモリ遅延を隠蔽しようとする戦略ではなく、メモリ階層の段数を減らすことで対応します。その結果として生じる NUMA の代償は許容範囲とします。
• 単一命令ストリームの ILP（命令レベル並列性）を高める：分岐予測や乱序実行コア、ハードウェア L1 キャッシュ、プリフェッチを活用し、Occupancy（同時に稼働するスレッドの割合）への依存を減らします。
• 適切な行列演算ユニットサイズを設定する：小規模なバッチでも性能が低下しないように設計します。
• Kernel の起動と同期に伴う固定オーバーヘッドを排除する。

推論 B：チップ面積はもはや希少な資源ではありません。むしろ、稼働していないマシン全体が浪費です。

同質クラスタを採用し、異種技術による分離を排除することで、専用リソースのスケジューリングボトルネックを防ぎます。

ピーク FLOPS を追求するのではなく、推論プロセスの各段階における実際の資源需要を十分に考慮し、計算能力・メモリ帯域幅・通信性能のバランスを保ちます。

一部の暗黙的なスロット（暗硅）を許容しつつ、チップが各段階の要件を満たすことを保証します。

チップ全体の消費電力を最適化し、これを設計上の主要目標と位置づけます。

通信遅延、特に長尾遅延（ロングテール・レイテンシ）を低減させます。

これらのアーキテクチャに関する意思決定は、以下の図に要約されています：

これらのアーキテクチャ上の考察は、OpenAI の考えとも一致しています。具体的な分析については次章で展開します。

1. OpenAI Jalapeño チップのアーキテクチャ

Jalapeño は、業界初の AI 向けハードウェアとソフトウェアを協調設計したプラットフォームです。ハードウェア面では、最初の RTL（レジスタ伝達レベル）記述から最終的な Tapeout（マスク作成）までわずか 9 ヶ月で完了しました。これはチップアーキテクチャの設計とも密接に関わっており、複雑な同期やハードウェアスケジューリング機構を排除したことで、検証期間が大幅に短縮されました。

また、市場投入までの期間短縮（タイム・トゥ・マーケット）のため、比較的成熟した Ethernet ベースの Scale-Up 技術を流用しました。メモリセマンティクスをサポートする標準 Ethernet を基盤とした業界初の Scale-Up は、私が 2020 年に開発した NetDAM が最初です。これは、実務的なトレードオフの結果でもあります。汎用スイッチチップはサプライチェーンが確立されておりコストも低く抑えられますが、遅延や有効利用率にはまだ課題が残ります。しかし、推論チップの相互接続に求められる極限の低遅延を実現するには、これは暫定的な解決策に過ぎません。後に Broadcom が発表した SUE にも大きな欠陥が見受けられました。同社と議論する際、私は FinePack のようなアプローチを参考にするよう提案しましたが、最終的に Broadcom は包括的な知識基盤を持っておらず、非常に詳細なキューイング遅延や、特に長尾遅延の抑制といった点を見落としていたことが判明しました。

3.1 チップアーキテクチャの概要

Jalapeño のチップアーキテクチャは以下の通りです：

imageこのチップは、Nvidia のような UMA（統一メモリアーキテクチャ）構造を放棄し、計算コアと HBM Slice を 1:1 で対応させることで、最小限の遅延でメモリアクセスを実現する経路を構築しました。さらに、チップ間ネットワーク（NoC）では、クロスダイ集合通信に特化した低遅延・高帯域幅の最適化が施されています。また、汎用イーサネットに基づくスケールアップ接続パスもバックアップとして残されています。

この設計はハードウェアとソフトウェアのコラボレーションの結果です。Jalapeño では多くのハードウェア機能が簡素化されており、統一された L2 キャッシュの削除や、カード間メモリセマンティクスの放棄、NUMA 距離の露出などが行われていますが、その複雑性はコンパイラやカーネル側に移されています。図の下部にある記述は、本質的に「明示的配置（explicit placement）」を意味しており、データがどのスライスに配置されるかが明確に決定され、コンパイラによって高度な最適化が行われます。この点については、後のソフトウェアシステムアーキテクチャの章で詳しく解説します。

また、このチップには中央调度器が存在せず、各コアが独立して命令を取得・実行します。これにより、「データ移動と同期」が実行時間の大部分を占めるという問題を回避しています。

本チップは複数のダイ（Die）をパッケージ化した構造を採用しており、その様子は以下の図の通りです：

image計算用ダイ（Compute Die）は TSMC の N3P 工艺で製造されており、リトレイルサイズ（reticle-size）の限界まで詰め込まれています。MXFP4 形式での総演算能力は 13.4 PFLOPS に達し、外部には最大 6 枚の HBM4 を接続可能です。

一方、I/O 機能は独立した IO-Die として分離され、TSMC の N3E 工艺が採用されています。計算と I/O を解耦リング（デカップリング）することで、次世代では計算用ダイのみを交換して IO チップレットを流用することも可能になり、その逆も可能です。これにより、Compute Die 上のリソース圧迫も軽減されます。立ち上げ間もないプロジェクトで迅速なイテレーションが求められる場合、こうした自由度は非常に価値があります。

なお、このチップのスケールアップ機能では、帯域幅を極限まで追求するのではなく、ローカル（ラック内）接続で 600GB/s、グローバル（ラック間）接続で 200GB/s をサポートしています。

また注目すべきは、Matrix エンジンが MXFP8 × MXFP4 の計算をネイティブにサポートしている点です。これは MoE（Mixture of Experts）モデルにおいて一般的に用いられる、活性化値 FP8 と重み FP4 の組み合わせパスに対応しています。さらに、Attention 計算については BF16 形式でのサポートも想定されています。

3.1.1 Core Slice アーキテクチャ

既知の公開情報に基づき、本チップの計算コア構造は以下のようになると推測されます：

imageまず、このプログラムは 1 つの GigaKernel を実行します。これは CUDA の Persistent Kernel に似ていますが、Jalapeño は遅延をより重視するため、CUDA Core のように複数の warp が占有されることはありません。代わりに、内部のスーパースケール構造によって構築された単一の命令ストリーム内のプリフェッチと乱序ウィンドウを利用して遅延を隠蔽します。これが、同プロセッサのフロントエンドにおける命令取得・デコード・乱序エミッションの仕組みです。

コア全体が高性能で動作し続けるためには、GPU の__syncthreads や mbarrier に相当する機能を実行できる、相対的に独立した同期アンカーポイントが必要です。

次に計算ユニットについてですが、通常は標準的な Tensor/Matrix Unit、SIMD/Vector Unit、Scalar Unit で構成されます。最も重要なのは Matrix Unit のアーキテクチャです。推論のデコードフェーズでは小バッチでも高い利用率を発揮し、全体としては Prefill といった計算集約型の処理も満たせることが求められます。ここがチップ設計において最も重要なポイントであり、どのように乱序実行と連携するか、最適な形状（Shape）は何か、遅延をどう隠蔽するかが鍵となります。これらについては後続のセクションで詳しく解説します。

続いてメモリサブシステムですが、Nvidia の SM が持つ RegisterFile / L1Cache / SMEM / TMEM といった階層構造とは異なり、L1 Cache のみを備えています。また、チップ全体にグローバルな L2 Cache も存在しません。HBM は各 Core Slice に直接スライス（Slice）として接続されており、これにより NOC 上の競合を効果的に回避しています。

Nvidia との最大の違いは、Tensor Unit と SIMD Unit / Scalar Unit の間でデータを SMEM / TMEM、Register File、L1 Cache の間を頻繁に転送・コピーする必要がない点です。両ユニットは L1 Cache を介して演算子を直接共有します。Tensor Unit が計算を終えた後、すぐに SIMD 命令を発行して継続して計算できます。これにより、バリアや DMA 編成といった固定コストが排除され、小規模な形状での計算もデコードフェーズの利用率に優しくなります。また、この L1 Cache は HBM に起因するジッター（揺らぎ）を効果的に吸収します。

3.1.2 NOC アーキテクチャ

NOC においては、Jalapeño が NUMA アーキテクチャを採用していることに加え、推論シーンで大量の並列操作における集合通信が必要となるため、OAI は集合通信を低遅延でサポートする独立した NOC を構築しました。同時に、NUMA ノード間の HBM メモリアクセスや ScaleUP IO Die の接続には、比較的低速な従来の NOC も併用しています。

3.2 Core Slice アーキテクチャの詳細解説

設計するこのチップは、Prefill（事前計算）、Decode（推論）、Draft（草案生成）、Verify（検証）の 4 つの処理段階において、いずれも業務ニーズをバランスよく満たすものでなければなりません。また、明確な弱点が生じないことと、性能対電力比（Perf-Watt Ratio）で数倍の向上を実現することも求められます。そこで、Decode 処理に焦点を当てて検討します。