GPU

概要

スループット、メモリ帯域、並列実行をつなげて理解する

GPU は「画面を描くための部品」から、現代では AI、HPC、映像処理、データ分析、科学技術計算、ブラウザ描画まで支える中核計算装置へ変わりました。このテキストでは、GPU を単なる高速なアクセラレータとしてではなく、CPU とは違う前提で設計された並列計算機 として理解します。

この章で重視すること

• GPU を「CPU の強化版」ではなく、設計目的の違う計算機として理解する
• ストリーミングマルチプロセッサ、warp / wavefront、共有メモリ、HBM、つまり高帯域メモリをひとつながりで捉える
• グラフィックス API と GPGPU、つまり GPU を描画以外の一般計算へ使う考え方を分けすぎず、同じハードウェアの別の顔として見る
• CUDA だけでなく ROCm / HIP、つまり AMD 系 GPU 向けの主要ソフトウェアスタックと移植層、Metal、Vulkan Compute まで視野に入れる
• 実務での判断材料として、向く処理・向かない処理・測るべき指標を押さえる

先に押さえる CPU と GPU の違い

GPU を理解しにくい最大の理由は、CPU の常識をそのまま持ち込んでしまうことです。CPU は Central Processing Unit の略で、OS、データベース、Web サーバ、アプリケーションの制御フローのような、分岐が多く、低遅延で順番に正しく処理したい仕事 を得意とします。GPU は Graphics Processing Unit の略で、よく似た計算を大量データへ一気に適用したい仕事 を得意とします。

一番大きな違い

• CPU は latency、つまり 1 個の仕事を終えるまでの待ち時間を小さくしたい
• GPU は throughput、つまり単位時間あたりの総処理量を大きくしたい

ここでいう latency は「1 件の処理が返るまでの速さ」、throughput は「全体としてどれだけ多く処理できるか」です。同じ「速い」でも、狙っている速さが違います。

CPU と GPU の比較表

先に知っておきたい専門用語

• メモリ帯域 は、1 秒あたりにどれだけ多くのデータを運べるかです。GPU はこの値が非常に重要です。
• レイテンシ隠蔽 は、あるスレッドがメモリ待ちの間に別のスレッドを動かして、待ち時間を表に出しにくくする考え方です。
• 並列性 は、同時に進められる仕事の量です。GPU は大きな並列性がある問題で真価を出します。
• ワークロード は、実際に処理したい仕事の性質です。GPU はワークロード適性の見極めがとても重要です。

CPU のコードをそのまま移しても速くならない理由

CPU 向けに自然なコードは、しばしば GPU 向けには不自然です。たとえば CPU では、

• 分岐を多用しても 1 スレッドが速ければ十分なことがある
• 小さなループを細かく関数分割しても問題になりにくい
• ポインタをたどる複雑なデータ構造でもキャッシュが効けば成立する

ということが起きます。GPU では逆に、

• 同じ warp 内で分岐が割れると足並みが崩れる
• 小さなカーネルを大量に起動すると起動オーバーヘッドが目立つ
• ランダムアクセスが増えると高い帯域を活かしにくい

となりやすいです。つまり GPU 化は、単なる「移植」ではなく、データ配置と計算形の再設計 になることが多いです。

CPU と GPU の見え方の違い

学び始めで持つとよい視点

GPU 学習の序盤では、次の順番で見ると理解しやすくなります。

1. 問題に十分な並列性があるか
2. データを連続的に読めるか
3. GPU 上へ乗せたデータを何度も再利用できるか
4. CPU との往復が支配的にならないか
5. 分岐や同期で実行束の足並みが崩れないか

目次

1. GPU とは何か
2. なぜ GPU は速いのか
3. GPU の歴史
4. 基本アーキテクチャ
5. グラフィックスパイプライン
6. GPGPU と CUDA 以後
7. 実行モデル：スレッド、warp、wavefront
8. メモリ階層
9. 同期、ダイバージェンス、占有率
10. テンソルコアと行列演算
11. GPU プログラミングモデル
12. 主要 API とエコシステム比較
13. CPU・TPU・NPU・FPGA との比較
14. 主要ユースケース
15. 性能最適化の勘所
16. マルチ GPU と相互接続
17. 仮想化、分割、マルチテナンシー
18. 運用・監視・トラブルシューティング
19. GPU が向かない処理
20. 2025-2026 の動向
21. 判断の指針
22. LLM 推論と GPU アーキテクチャ
23. GPU 最適化の実践パターン
24. CUDA・ROCm・Metal・Vulkan のコードモデル比較
25. FAQ
26. 参考文献

GPU とは何か

GPU は Graphics Processing Unit の略で、もともとは画像描画を高速化するために発展したプロセッサです。ただし現代では、単に「グラフィックス専用」ではありません。むしろ、

• 同じ種類の計算を大量のデータへ適用する
• 非常に高い メモリ帯域、つまり大量データを高速に読み書きする力が必要
• 単一スレッドの速さより全体の 処理量 が重要

という領域で、CPU よりはるかに強い計算機として使われています。

GPU を一言でいうと

GPU は、少数の重いコアで低遅延処理を得意とする CPU に対し、多数の軽量な実行資源で高スループット処理を得意とする装置 です。ここでいう コア は、命令を実行する計算の中心部分です。ただし GPU の「コア」は CPU コアと同じ意味ではなく、より細かい実行資源の集合として扱う方が実務では正確です。

GPU を理解すると何が見えるか

GPU を理解すると、

• なぜ AI 学習で GPU が主役なのか
• なぜゲームや 3D 描画が GPU 依存なのか
• なぜ CPU のコードをそのまま移しても速くならないのか
• なぜメモリ転送やダイバージェンスがボトルネックになるのか

が現実的に見えてきます。

GPU を誤解しやすいポイント

• コア数が多い = いつでも速い ではない
• VRAM が多い = いつでも高性能 ではない
• 占有率が高い = いつでも最速 ではない
• GPU 利用率が 100% に近い = 最適化済み でもない

これらはすべて一部しか見ていません。GPU では、演算器、メモリ帯域、転送、同期、ソフトウェアスタックが互いに制約になります。

なぜ GPU は速いのか

GPU が速い理由は、クロック周波数、つまり 1 秒あたりの動作回数が極端に高いからではありません。主な理由は次の 4 つです。

1. 実行資源の数が多い
2. メモリ帯域が非常に広い
3. レイテンシを隠す設計になっている
4. 同じ命令を大量データへ適用する処理に強い

レイテンシよりスループット

CPU は「1 本の重要な仕事をできるだけ早く終わらせる」方向に強く最適化されています。GPU は「大量の似た仕事をまとめてさばく」方向に最適化されています。CPU は分岐予測、投機実行、大きなキャッシュを使って 1 スレッドの待ち時間を減らします。GPU は多数のスレッドを切り替えながら動かして、待ち時間そのものを別の仕事で埋めます。

GPU の速さは万能ではない

GPU は万能高速化装置ではありません。次のような処理は苦手です。

• 分岐が多く、各要素で違う経路を通る
• データ依存が強く、並列化しにくい
• ランダムアクセスが多い
• GPU に送る前後の転送コストが支配的

GPU の速さは、向いた形に問題を変形できたときに出る と考える方が正確です。

GPU の歴史

固定機能時代

初期の GPU は、現在のような汎用計算機ではありませんでした。テクスチャマッピング は画像を物体表面へ貼る処理、ラスタライズ は三角形などの図形を最終的な画素へ変換する処理、Z バッファ は前後関係を判定するための深度情報、ブレンディング は色を混ぜ合わせる処理です。初期 GPU はこうしたグラフィックス専用の固定機能を高速に実行する装置でした。

プログラマブルシェーダの登場

2000 年代に、頂点シェーダ、つまり 3D モデルの頂点ごとに実行するプログラムと、ピクセルシェーダ、つまり画素ごとに色や光を計算するプログラムがプログラマブルになり、描画処理の一部をソフトウェア的に制御できるようになりました。ここで「GPU 上で計算する」発想が一気に強まりました。

GPGPU

当初の GPGPU は General-Purpose computing on GPUs の略で、GPU を描画以外の一般計算へ使うことです。初期の GPGPU は、グラフィックス API を無理に計算へ転用する形でした。行列をテクスチャとして扱ったり、レンダリングパスを計算として流用したりしていたため、かなり不自然でした。

CUDA と汎用化

2006 年に NVIDIA が CUDA を出したことで、GPU を C/C++ 風に直接プログラムできる道が開きました。これが、現在の GPU 計算の大きな分岐点です。

現代

現在の GPU は、

• グラフィックス
• 汎用計算
• 行列演算
• レイトレーシング
• ビデオエンコード / デコード

など複数の用途を 1 チップへ統合しています。

基本アーキテクチャ

GPU の細部はベンダごとに違いますが、概念的にはかなり似ています。

• 実行クラスタが複数ある
• 各クラスタが多数スレッドを切り替えながら処理する
• 高帯域なメモリへ接続される
• 命令キャッシュ、L1、つまり実行クラスタの近くにある小さなキャッシュ、共有メモリ、L2、つまり GPU 全体で共有するより大きなキャッシュを持つ

NVIDIA 系の見方

NVIDIA では主に SM(Streaming Multiprocessor)という単位で GPU を見ます。SM は、多数のスレッドをまとめて走らせる実行クラスタです。実際の命令発行器、レジスタファイル、共有メモリ、各種演算器がこの単位の周辺にまとまっています。

AMD 系の見方

AMD では主に CU(Compute Unit)や、最近では WGP(Work Group Processor)といった単位が重要です。名前は違いますが、NVIDIA の SM と同様に、多数の並列スレッドを処理する実行クラスタとして捉えると理解しやすくなります。

Apple / Metal 系の見方

Apple は API 面では Metal、つまり Apple の低レベル GPU API を前面に出し、GPU family と feature set で能力を表します。内部構造は公開範囲が少し異なりますが、実務的には スレッドグループ、つまり共同で共有メモリを使う実行単位と、タイルベース最適化、つまり画面を小さな領域に分けて効率よく処理する考え方が重要です。

GPU で重要な構成要素

• 実行クラスタ: SM や CU のように実際の並列実行を担うまとまり
• レジスタファイル: 各スレッドが最も近い場所で使う超高速な小容量メモリ
• 共有メモリ / LDS / threadgroup memory: 同じブロックやグループ内のスレッドが共有できる手動管理メモリ。LDS は AMD 系でよく使う Local Data Share の呼び名
• L1 / texture cache: 近くで再利用されるデータを貯める小さな高速キャッシュ
• L2: GPU 全体で共有される比較的大きなキャッシュ
• グローバルメモリ: VRAM、つまり GPU ボード上の主記憶
• コピーエンジン: CPU メモリと GPU メモリの転送などを担当する専用経路
• ビデオ専用エンジン: エンコードやデコードを担当する専用回路
• テンソル / 行列演算ユニット: 行列積や AI 演算向けに特化した専用演算器

CPU と GPU の内部構造の違い

CPU は、複雑な制御フローや低遅延応答のために、分岐予測器、投機実行、大きなキャッシュ、アウトオブオーダ実行を多く抱えます。アウトオブオーダ実行 は、プログラム順ではなく、依存関係がない命令を先に進めて待ち時間を減らす仕組みです。

GPU は、そうした複雑な制御機構へ面積と電力を大量に使う代わりに、

• 多数の実行レーン
• 大きなレジスタファイル
• 高帯域メモリ接続
• 多数スレッドの高速切替

へ資源を振ります。これは「1 本を速く」より「束全体を多く流す」ための選択です。

GPU を制御プレーンとデータプレーンで見る

• 制御プレーン: CPU、ドライバ、ランタイム、コマンドキュー、同期 API
• データプレーン: 実際に GPU 上で動くカーネル、メモリ転送、共有メモリ、演算器

実務では、性能が悪いときに「GPU カーネルの中」だけを見がちですが、実際には CPU 側のデータ供給やランタイムの同期がボトルネックなことも多いです。

グラフィックスパイプライン

GPU は今もグラフィックス用の装置でもあります。API は Application Programming Interface の略で、ソフトウェアから GPU 機能を呼び出すための約束事です。グラフィックス API を理解すると、GPU の固定機能とプログラマブル部分の境界が見えます。

なぜこの章が重要か

GPU を AI 用アクセラレータとしてしか見ないと、もともとの設計思想を見失いやすいです。テクスチャキャッシュ、補間、タイル、レンダーパス最適化などは、いまでも GPU の物理設計に影響を残しています。

Compute との関係

現代の GPU は、同じハードウェア資源を

• グラフィックス
• コンピュート
• コピー
• メディア処理

で共有します。だから、グラフィックス API とコンピュート API は別世界ではありません。

グラフィックス由来の設計が compute に残す影響

GPU はもともと描画のために進化したため、compute だけを見ると見落としやすい性質が残っています。

• テクスチャキャッシュは空間局所性の高いアクセスで効きやすい
• タイルベース設計は一時データを近いメモリへ閉じ込めやすい
• ラスタライズや補間のための固定機能は、専用回路の価値を示す好例でもある

この視点を持つと、なぜ AI 時代にも「専用ユニット」が増えるのかが理解しやすくなります。

GPGPU と CUDA 以後

GPGPU の本質

GPGPU は、General-Purpose computing on GPUs の略です。つまり GPU を、描画以外の一般計算へ使うことです。

CUDA の意味

CUDA は NVIDIA の GPU 計算プラットフォームです。重要性は、NVIDIA 専用という点だけではありません。GPU 計算を

• カーネル起動
• スレッド階層
• メモリ階層
• 同期
• ホストとデバイスの分離

という形で、比較的一貫したプログラミングモデルへ落とし込んだことにあります。

現代の主要系統

• CUDA: NVIDIA の GPU 向けに最も広く使われる計算基盤
• HIP / ROCm: AMD 系 GPU 向けの主要スタック。HIP は CUDA に近い書き味の移植層
• Metal: Apple プラットフォームの GPU API
• Vulkan Compute: Khronos の低レベルクロスベンダ API の compute 機能
• OpenCL: 標準系だが、近年の AI 実務の中心ではない
• Direct3D 12 Compute: Windows 系で使われる低レベル GPU compute

2026 年の実務感覚では、この並びは単なる API 一覧ではありません。ざっくり言うと、

• 学習・推論・周辺ライブラリの厚み では CUDA が依然として非常に強い
• HPC / 移植性 / オープン寄りの基盤 では ROCm の重要性が上がっている
• Apple 統合環境 では Metal がほぼ前提になる
• ブラウザ / サンドボックス / 配布容易性 では WebGPU 系の価値が目立つ

という棲み分けが見えます。

なぜ CUDA が強いのか

CUDA の優位は API だけではありません。実務上は、

• コンパイラやランタイムが成熟している
• cuBLAS や cuDNN のような基盤ライブラリが厚い
• NCCL のようなマルチ GPU 通信ライブラリが強い
• プロファイラやデバッガが揃っている
• 上位フレームワークの最適化が先に来やすい

という、ソフトウェア生態系の積み上げが非常に大きいです。

2026 年時点でも、この「エコシステム差」はかなり効きます。GPU の理論性能が近くても、

• 使いたいライブラリがすぐ動くか
• profiler や debugger が十分に揃っているか
• 推論ランタイムや量子化ツールが枯れているか
• 周辺フレームワークの最適化が先に来るか

で、現場の生産性は大きく変わります。

ROCm / HIP の価値

ROCm は単に「CUDA の代替」ではありません。AMD 系 GPU を前提にしつつ、HIP、ライブラリ群、プロファイラ、フレームワーク連携を含んだ総合基盤です。複数ベンダ調達やコスト、オープン性、HPC 既存基盤との整合を重視する組織では、十分に現実的な選択肢です。

しかも 2026 年時点の ROCm 公式 docs は、ROCm を HIP / OpenCL / OpenMP を含む統合ソフトウェア基盤 として整理しています。ここはかなり大事です。つまり AMD 系では、

• CUDA 風のカーネル移植なら HIP
• 既存 HPC コードの GPU offload なら OpenMP offloading
• 互換性や既存資産との接続では OpenCL

のように、入口が 1 つではありません。ROCm は「AMD 上の CUDA もどき」だけで見ると、少し狭く理解しすぎになります。

実行モデル：スレッド、warp、wavefront

GPU プログラミングで最初に混乱しやすいのが、CPU の スレッド と GPU の スレッド は同じ重さではない、という点です。CPU のスレッドは OS スケジューラが扱う比較的重い実行単位ですが、GPU のスレッドはもっと軽く、大量に作ってまとめて処理する前提で設計されています。

CUDA の基本階層

warp

NVIDIA では、通常 32 スレッド単位で命令が進みます。これを warp と呼びます。warp は「GPU が内部でまとめて進めるスレッドの束」です。

wavefront

AMD では、同様の概念として wavefront が重要です。世代やモードにより幅の扱いは異なりますが、CPU の 1 スレッドとはかなり違う粒度で動いています。実務では「warp に近い実行束」と考えると概ね外しません。

SIMT

GPU はしばしば SIMT と説明されます。これは Single Instruction, Multiple Threads の略で、見た目は多数スレッドですが、実際には同じ命令流を束で実行する側面が強い、という考え方です。SIMD が 1 命令で複数データを処理するベクトル計算だとすると、SIMT はそれを「多数スレッドがあるように見せる」形へ寄せたものです。

ダイバージェンス

同じ warp の中で条件分岐により別の経路へ進むと、GPU はしばしばそれらを逐次化して扱います。これが warp divergence、つまり分岐による足並み崩れです。

kernel / grid / block / threadgroup / subgroup

GPU の用語は API ごとに少しずつ違うので、対応関係を早めに押さえると読みやすくなります。

GPU スケジューリングの感覚

GPU は大量スレッドを同時に「本当に全部動かす」わけではありません。実際には、SM や CU ごとに載せられる warp / wavefront 数には上限があり、

1. まずカーネルが起動される
2. スレッドがブロックや workgroup にまとめられる
3. それらが各実行クラスタへ割り当てられる
4. 実行可能な warp / wavefront から順に切り替えながら進む

という流れになります。ここでレジスタや共有メモリを使いすぎると、同時に載せられる束の数が減り、レイテンシ隠蔽力が落ちます。

async compute

async compute は、描画や別キューの処理と並行して compute を流す考え方です。ただし「必ず速い」わけではありません。共有するメモリ帯域やキャッシュの競合、同期の増加で逆効果になることもあります。

メモリ階層

GPU の性能は、かなりの割合でメモリで決まります。CPU 以上に「どこから、どの順番で、どれだけまとめてデータを読むか」が効きます。

GPU の主なメモリ層

• レジスタ
• 共有メモリ / LDS / threadgroup memory
• L1 / read-only cache / texture cache
• L2
• グローバルメモリ
• ホストメモリ

共有メモリ

共有メモリは、スレッドブロック内の協調で極めて重要です。低レイテンシですが容量は小さく、使いすぎると 占有率 に影響します。占有率は、1 つの SM や CU に何組の実行束を同時に載せられるかの目安です。

グローバルメモリ

容量は大きいですが遅いです。GPU の最適化では、グローバルメモリアクセスの回数とパターンが性能を大きく左右します。

コアレッシング

隣接スレッドが隣接アドレスを読むと、ハードウェアがまとめて効率よく転送しやすくなります。これが coalesced access、日本語ではアクセスのまとめ読みです。CUDA の Best Practices Guide でも、グローバルメモリアクセスのコアレス化は高優先度の最適化として強調されています。

HBM

AI / HPC 向け GPU では HBM、つまり High Bandwidth Memory が重要です。HBM は非常に高い帯域を実現しますが、容量やコスト、実装複雑性とのトレードオフもあります。

register pressure

register pressure は、1 スレッドあたりのレジスタ使用量が多すぎて、同時実行数やコンパイラ最適化を圧迫する状態です。GPU 最適化では「メモリへ出さないためにレジスタを使う」ことと、「使いすぎて占有率を落とす」ことのバランスが常にあります。

bank conflict

共有メモリは速いですが、アクセスの仕方によっては bank conflict、つまり同時アクセスの衝突が起きます。これは、複数スレッドが内部的に同じバンクへ集中的にアクセスし、直列化が増える現象です。共有メモリは「使えば速い」ではなく、「並び方まで設計して速い」です。

pinned memory と unified memory

• pinned memory はページ固定されたホストメモリで、DMA 転送を効率化しやすい
• unified memory は CPU と GPU から見えるメモリ空間を統合して扱いやすくする仕組み

どちらも便利ですが、便利さと最高性能は同じではありません。とくに unified memory はアクセスパターンによって暗黙の移動が起きるため、性能の説明が難しくなることがあります。

良いアクセスと悪いアクセス

同期、ダイバージェンス、占有率

同期

GPU は大規模並列ですが、何でも自由に同期できるわけではありません。まずは

• ブロック内同期
• グリッド全体同期
• ホスト側同期

の違いを分ける必要があります。

占有率

occupancy は、ある SM / CU に対してどれだけ多くの warp / wavefront を同時に抱えられるか、という概念です。一般に高すぎても低すぎても単純な正義ではありませんが、低すぎるとレイテンシ隠蔽が難しくなります。

占有率に効くもの

• レジスタ使用量
• 共有メモリ使用量
• ブロックサイズ
• ハードウェア上限

占有率は手段であって目的ではない

占有率を上げるためにレジスタ圧縮やブロックサイズ調整を行うことはありますが、最終的に見るべきなのは

• 実行時間
• スループット
• メモリ帯域利用率
• 演算ユニット利用率

です。

同期で詰まりやすいポイント

• ブロック内同期を多用しすぎる
• 原子操作にホットスポットが集中する
• グローバル同期が必要なアルゴリズムをそのまま持ち込む
• CPU 側の同期で GPU を頻繁に待たせる

原子操作 は、複数スレッドが同時に更新しても壊れないようにする操作です。便利ですが、同じ場所へ多数スレッドが集まると強い競合が起きます。

ダイバージェンスを避ける考え方

• 条件分岐をデータ前処理で減らせないか
• 同じ分岐を通る要素同士を近くへ並べられないか
• マスク計算や別カーネル分割の方が有利ではないか

GPU では「分岐をゼロにする」ことより、「同じ実行束の中で分岐を揃える」ことが大事です。

テンソルコアと行列演算

AI 時代の GPU を語るうえで、テンソルコア や行列エンジンは外せません。テンソルは多次元配列、テンソルコアはその配列演算、とくに行列積を高密度に処理する専用回路です。

何が違うのか

通常の ALU、つまり加減算などの基本演算器や、FP、つまり浮動小数点演算ユニットとは別に、行列積や fused multiply-add、つまり掛け算と足し算をまとめて行う演算を高密度に処理する専用ユニットを持つことで、深層学習や線形代数に対して圧倒的なスループットを出せます。

なぜ AI に効くのか

深層学習の多くは、結局は大規模な行列演算、畳み込み、テンソル演算へ帰着します。そこに専用ユニットを当てることで、通常演算器だけよりはるかに効率が上がります。

ただし注意点

• 精度形式の選択が重要
• メモリ転送が詰まるとユニットが遊ぶ
• 小さすぎる問題では効かない

学習と推論で見え方が違う

• 学習 は前向き計算、逆伝播、勾配更新を含むため、計算量もメモリ使用量も大きい
• 推論 は学習より軽いことが多いが、レイテンシ、バッチサイズ、キャッシュ再利用が重要になる

とくに LLM 推論では、行列演算そのものだけでなく、KV cache、つまり過去トークンのキーとバリューを保持するメモリ構造が性能と容量の大きな制約になります。

GPU プログラミングモデル

CUDA

CUDA は、ホストコードからカーネルを起動し、GPU 上で大量スレッドを動かすモデルです。

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

このコード自体は単純ですが、実際には

• メモリ配置
• ブロックサイズ
• 転送コスト
• 同期

が性能を左右します。

HIP / ROCm

HIP は CUDA に近い書き味を持ちつつ、AMD GPU を主対象にしたプログラミングモデルです。移植性を重視する現場ではかなり重要です。ROCm の公式資料でも、HIP、OpenCL、OpenMP を含む総合ソフトウェアスタックとして説明されています。

Metal

Apple の Metal は、グラフィックスとコンピュートを一貫した API で扱います。Apple GPU の feature family を前提に、スレッドグループ単位の設計が重要です。

Metal を実務で見るときは、「Apple GPU 向けの compute API」というだけでなく、

• OS / ドライバ / GPU / メモリモデルを Apple が一体で設計している
• graphics と compute を同じパイプライン設計思想で扱える
• unified memory 前提で CPU / GPU のデータ配置設計が変わる

という点が大きいです。Apple 系では、単に CUDA がない代替 API ではなく、Apple 向けに自然な GPU の入口 として理解した方が噛み合います。

Vulkan Compute

Vulkan は明示的な低レベル API で、グラフィックスだけでなく compute も扱えます。抽象化は少なめですが、制御の自由度が高いです。Khronos の Vulkan Guide でも、Vulkan は graphics と compute の両方を必須機能として備える API と整理されています。

ランタイムで起きていること

どの API を使っていても、実際には次のような流れが起きています。

1. CPU 側でバッファやテクスチャなどの資源を用意する
2. GPU 用のカーネルやシェーダをコンパイル、またはロードする
3. コマンドバッファやストリームへ処理を積む
4. GPU がそれを非同期に実行する
5. 必要なら結果を CPU 側へ戻す

GPU プログラミングの難しさは、この「CPU の時間」と「GPU の時間」がずれる点にあります。見かけ上は 1 行の API 呼び出しでも、裏では非同期に長い仕事が走ることが珍しくありません。

ROCm 導入で先に確認したいこと

ROCm を本番候補として見るときは、性能を測る前に次を確認した方が安全です。

1. 使う GPU が、その ROCm バージョンで正式サポートされているか
2. Linux / Windows、カーネル、ドライバ、コンテナ環境の組み合わせが対応範囲か
3. 使いたいフレームワーク機能が HIP / OpenMP / OpenCL のどこで乗るのか
4. profiler や debugger を含めた運用ツールが現場の要件を満たすか

ROCm は「動くかどうか」だけでなく、どのレイヤで使うつもりか を先に決めると判断しやすくなります。

主要 API とエコシステム比較

実務上の見方

• NVIDIA 中心なら CUDA が最短
• AMD を含めるなら HIP を真剣に検討
• Apple プラットフォームなら Metal が基本
• 低レベルで長寿命なクロスプラットフォーム基盤なら Vulkan Compute を候補に入れる

API 選定を 1 枚で見る

API 選定フロー

もう一歩実務寄りに言うと、WebGPU は「性能の極限」よりも、

• ブラウザでそのまま配れる
• セキュアなサンドボックスで GPU を使える
• OS や GPU ベンダをまたぎやすい

ことに価値があります。ネイティブ GPU API の完全な代替ではありませんが、配布しやすい GPU 計算 という意味ではかなり強いです。

主要ライブラリと周辺ツール

API を選ぶときの実務的な問い

• どの GPU ベンダを主対象にするのか
• 手元で最速にしたいのか、長期の移植性を取りたいのか
• 上位ライブラリの恩恵を強く受けるのか、独自カーネルを書くのか
• AI 推論や HPC のように、既存エコシステムが結果を大きく左右する分野か

WebGPU / WGSL をどう位置づけるか

WebGPU を誤解しやすいのは、「ブラウザ版 CUDA」だと思ってしまうところです。実際には、

• ブラウザで GPU を使える
• 安全な実行モデルを重視している
• WGSL という専用言語で、表現力と移植性のバランスを取っている

という設計です。

実務で向いているのは、

• 画像処理や可視化のデモ
• 学習教材
• その場で試せる軽量推論
• クライアントサイド処理の高速化

であって、巨大モデルの本格学習や極限チューニングの中心ではありません。逆に言えば、配布のしやすさと再現性 が価値になる場面ではかなり強いです。

どのレイヤで戦うか

CPU・TPU・NPU・FPGA との比較

GPU の立ち位置

GPU は、汎用性と性能のバランスが非常によい高速並列装置です。だからこそ広く使われています。CPU ほど自由ではないが、TPU、つまり行列演算へかなり特化した AI アクセラレータや、NPU、つまり端末や SoC に載る AI 専用演算器、FPGA、つまり回路自体を再構成できるデバイスほど用途を狭めなくてもよい、という中間の強さがあります。

主要ユースケース

3D レンダリング

古典的な GPU の用途です。頂点処理、ラスタライズ、シェーディング、ポストプロセスなどを高速に行います。

深層学習

学習でも推論でも、GPU は中心的です。特に大規模行列演算と高帯域メモリが効きます。

学習では、巨大なミニバッチや活性値、勾配、オプティマイザ状態を持つため、計算能力だけでなくメモリ容量と通信が効きます。推論では、単発リクエストのレイテンシを詰めたい場面と、多数リクエストを束ねてスループットを稼ぎたい場面で最適解が変わります。

HPC

気象、分子動力学、有限要素法、数値流体、天体計算など、大規模数値計算に強いです。

映像処理

エンコード、デコード、画像フィルタ、超解像、映像推論などで GPU は非常に有効です。

データ分析

列指向処理、ベクトル演算、大規模フィルタリングや group-by の一部を GPU へ逃がすケースも増えています。

ただし、GPU が効きやすいのは列指向で連続アクセスしやすい処理です。複雑な結合、文字列処理、分岐だらけの UDF が支配的だと、期待ほど伸びないこともあります。

ブラウザと UI

CSS、Canvas、WebGL / WebGPU、動画再生など、多くの UI はすでに GPU 依存です。

ブラウザ文脈では、専用 GPU だけでなく統合 GPU やモバイル GPU も重要です。ここでは純粋な最大性能より、電力効率、タイルベース最適化、メモリ共有モデルが効きます。

LLM 推論

現代の GPU ユースケースとして特に重要なのが LLM 推論です。ここでは、

• 重みを VRAM に載せ切れるか
• KV cache をどれだけ効率よく保持できるか
• リクエストをバッチングできるか
• 量子化しても品質が許容されるか

が支配的です。単に FLOPS が高いだけでは足りず、VRAM 容量、帯域、通信、サービングランタイムの賢さが性能を大きく左右します。

ユースケース別に見る GPU の論点

性能最適化の勘所

GPU 最適化でまず疑うべきは次の順です。CUDA Best Practices Guide でも、まずプロファイルを取り、実際のホットスポットとボトルネックを見つけることが高優先度で推奨されています。

1. ホストとデバイスの転送が支配していないか
2. グローバルメモリアクセスが非効率ではないか
3. ダイバージェンスが激しくないか
4. ブロックサイズや threadgroup サイズが極端ではないか
5. 小さすぎる問題を GPU へ投げていないか

典型的なボトルネック

• PCIe、つまり CPU と GPU をつなぐ汎用高速バスの転送
• メモリ帯域不足
• L2 / shared memory の使い方
• occupancy 低下
• launch overhead

プロファイリング

GPU では、推測より計測が重要です。一般に見るべき指標は、

• カーネル実行時間
• メモリ転送時間
• SM / CU 利用率
• メモリ帯域利用率
• warp stall 理由

です。

compute bound と memory bound

GPU カーネルは、大雑把には次の 2 つに分けて考えると整理しやすいです。

• compute bound: 演算器が支配的で、もっと計算器が速ければ伸びる
• memory bound: データ供給が支配的で、帯域やアクセス効率が悪い

この切り分けを間違えると、無意味な最適化へ時間を使いやすくなります。メモリ律速なのに演算を少し減らしても効果は薄く、逆に計算律速なのにアクセス最適化ばかりしても伸びません。

roofline 的な見方

roofline model は、演算強度、つまり「1 byte あたりどれだけ計算するか」と、到達できる性能の上限を結びつけて考える枠組みです。細部は高度ですが、実務的には「そのカーネルは計算を増やしてもよいのか、まずデータ再利用を増やすべきか」を整理するのに役立ちます。

最適化の実務チェックリスト

1. まず end-to-end で時間を測る
2. CPU 時間、転送時間、GPU カーネル時間を分ける
3. 最も重いカーネル 1 つに絞る
4. そのカーネルが帯域律速か計算律速かを判定する
5. アクセス、分岐、占有率、同期の順に疑う
6. 改善後に必ず再計測する

この順番は、NVIDIA の CUDA Best Practices Guide の考え方ともかなり一致しています。要するに GPU 最適化は、

1. Assess: どこが重いかを測る
2. Parallelize: 並列化できる部分を見極める
3. Optimize: いちばん重い箇所だけを詰める
4. Deploy: 改善後の全体効果を確認する

という循環で進めた方がうまくいきます。最初から occupancy や shared memory の細部へ飛び込むより、まずホットスポットの位置を測る方がずっと強いです。

典型的な改善パターン

• 小さいカーネルをまとめる
• データ転送回数を減らす
• タイリングして共有メモリ再利用を増やす
• ブロックサイズを調整する
• 混合精度や量子化を使う
• CPU 側前処理を見直して分岐を減らす

マルチ GPU と相互接続

1 枚の GPU では足りない場合、複数 GPU を束ねます。

接続

• PCIe
• NVLink、つまり NVIDIA の GPU 間高帯域相互接続
• Infinity Fabric 系
• CPU ソケット越し

何が難しいか

マルチ GPU では、

• データ分割
• 通信
• 同期
• 負荷偏り

が難しくなります。単純に 2 倍、4 倍にはなりません。

データ並列とモデル並列

AI では、

• データ並列
• テンソル並列
• パイプライン並列

などの並列化手法が使われます。

通信が支配的になる瞬間

マルチ GPU では、単体 GPU の最適化とは別の世界が始まります。とくに学習では、勾配同期の all-reduce が支配的になりやすく、推論ではモデル並列の境界で待ちが増えます。ここでは GPU 自体の速さだけでなく、

• GPU 間リンクの帯域
• トポロジ
• 通信ライブラリ
• バッチ設計

が結果を左右します。

仮想化、分割、マルチテナンシー

GPU は高価なので、1 人 1 枚専有ではもったいない場面が多いです。

MIG

NVIDIA の MIG は Multi-Instance GPU の略で、1 枚の物理 GPU を複数の独立インスタンスへ分割する仕組みです。MIG User Guide の対応表では、Ampere 世代以降で対応が始まり、2026 年 4 月時点では Hopper や Blackwell 系の対応 GPU も並んでいます。

さらに 2026 年時点の NVIDIA 公式ガイドでは、Hopper 世代以降では MIG 有効化時に GPU reset が不要 という挙動差も明記されています。ここは運用上かなり重要です。Ampere 世代の感覚で「有効化は重い操作」と思い込むと、世代差を見誤ります。

もうひとつ実務で見落としやすいのは、MIG モードの持続性がドライバ常駐に依存する ことです。つまり「一度設定すれば永続」と思い込むと、再起動やドライバ再読込時の自動化設計で詰まりやすくなります。

何が嬉しいか

• マルチテナント運用しやすい
• QoS を読みやすい
• 小さなジョブを詰めやすい

何が難しいか

• 分割設計
• 再起動や再構成運用
• Kubernetes とデバイスプラグイン連携
• 監視粒度

MIG と time slicing の違い

MIG はハードウェア資源を分割して、比較的強い分離を与える仕組みです。一方 time slicing は、時間で順番に GPU を使わせる共有です。

小さい推論ジョブを多数収容したいのか、厳密な分離が必要なのかで選び方が変わります。

運用・監視・トラブルシューティング

まず見るもの

• GPU 利用率
• メモリ使用量
• 温度と電力
• 転送帯域
• エラー
• ドライバ / ランタイム版

典型的なトラブル

1. GPU 利用率が低い

• CPU 側で詰まっている
• バッチサイズが小さい
• 転送待ちが長い
• カーネルが小さすぎる

ここでいう GPU 利用率は 1 指標にすぎません。利用率が低くても、短いバースト処理なら正常なことがありますし、逆に利用率が高くてもメモリ待ちだらけで効率が悪いこともあります。

2. メモリ不足

• モデルやバッチが大きすぎる
• 一時バッファが多い
• フラグメント化
• mixed precision を使っていない

3. 速くならない

• GPU 向きでない処理
• 実は I/O が支配的
• カーネル launch が細かすぎる
• ダイバージェンスが強い

4. 再現性が揺れる

• 並列 reduction の順序差
• 非決定的アルゴリズム
• 原子操作順序

監視で見たいメトリクス

• GPU 使用率
• SM / CU 使用率
• メモリ使用量と残量
• メモリ帯域利用率
• PCIe / NVLink 転送
• 温度、消費電力、クロック低下
• ECC や Xid などのハードウェアエラー
• キュー長、推論待ち時間、バッチサイズ

切り分けの順番

1. まず end-to-end の遅さか、GPU 単体の遅さかを分ける
2. CPU / I/O / 転送が支配していないか確認する
3. GPU 内では最も重いカーネルを特定する
4. メモリ律速か計算律速かを判断する
5. 再現性の揺れやドライバ差分を確認する

GPU が向かない処理

GPU を使えば常に速いわけではありません。典型的に向きにくいのは次です。

• 単発で短い処理
• ポインタ追跡が多い
• 強い逐次依存
• 分岐だらけのロジック
• データ量が少ない
• CPU との往復が多い

判断の基本

「同じ処理を大量データへ当てるか」「データを GPU 上にしばらく留められるか」をまず見ると、大きく外しにくいです。

GPU を使わない方がよい具体例

• 1 回しか呼ばない小さな前処理
• 複雑な木構造やグラフをポインタでたどる処理
• 1 件ごとに強い分岐が入るルールエンジン
• リクエストごとにデータ転送だけで終わるような処理

こうしたケースでは、CPU でまず速く作る方が全体最適になることがよくあります。

2025-2026 の動向

AI 中心の設計が続く

GPU は引き続き AI インフラの中心です。行列演算性能、HBM 容量、相互接続、ソフトウェアスタックが差別化点になっています。

GPU 分割と共有の重要性が上がる

高価な GPU を効率よく使うため、MIG や Kubernetes 連携、ワークロード分離の重要性がさらに上がっています。

CUDA と ROCm の二極化が濃くなる

NVIDIA 中心の CUDA 優位は強いままですが、AMD ROCm / HIP も HPC と一部 AI ワークロードで存在感を増しています。2026-03-10 更新の ROCm 公式 docs でも、HIP、OpenCL、OpenMP を含む統合スタックとして整理されています。

Apple は Metal / Unified Memory を一体で進化

Apple の GPU は、Metal feature family と unified memory 前提の設計が重要です。CPU / GPU / NPU の境界をまたぐ設計判断が特徴的です。

WebGPU / Vulkan Compute の裾野拡大

ブラウザやクロスプラットフォーム compute の文脈で、WebGPU や Vulkan Compute の重要性も増しています。低レベル API への理解が、グラフィックスだけでなく汎用計算にも効きます。

WebGPU の現在地を雑に言うと、「ブラウザで GPU をまじめに使うための標準の入口」 です。WGSL という専用シェーディング言語と組み合わせて、

• 画像処理
• 軽量な機械学習推論
• 可視化
• 教材・デモ

のような用途でかなり扱いやすくなっています。ネイティブ API の完全な代替ではありませんが、「配布しやすい GPU 計算」という意味ではとても重要です。

サービングランタイムの重要性が上がる

2025-2026 の実務では、GPU ハードウェア単体よりも、推論サーバ、バッチング戦略、量子化、キャッシュ管理、マルチテナント制御の差で結果が大きく変わる場面が増えています。とくに LLM では、GPU を「どう使うか」のソフトウェア層が性能の大部分を決めることも珍しくありません。

判断の指針

GPU を使うべき場面

• 行列演算やテンソル演算が中心
• 大量データへ同じ処理を適用
• バッチ処理でまとめられる
• 高メモリ帯域が必要
• レイテンシよりスループットが重要

CPU のままの方がいい場面

• 制御分岐中心
• 小規模処理
• 逐次依存が強い
• 転送が支配的

まず小さく検証する

実務では、最初から GPU 前提で巨大設計にするより、

1. CPU 実装のボトルネックを確認
2. GPU 化候補を分離
3. 小さくベンチマーク
4. 転送コスト込みで判断

の順が安全です。

導入判断の 5 つの質問

1. この処理には大量並列にできる部分が本当にあるか
2. データは GPU 上へしばらく置いて再利用できるか
3. 既存ライブラリで大部分を賄えるか
4. 運用で必要な監視、分割、共有方式を用意できるか
5. GPU コストに見合うスループット改善があるか

LLM 推論と GPU アーキテクチャ

LLM 推論は、現代の GPU を理解するうえで最も重要な実例のひとつです。ここでは「なぜ LLM が GPU を強く必要とするのか」と「どこで詰まるのか」を、ハードウェアとランタイムの両面から整理します。

推論の大まかな流れ

LLM 推論の時間感覚

prefill と decode

• prefill は、最初に与えられた文脈全体を処理して内部状態を作る段階です
• decode は、その後 1 トークンずつ自己回帰的に続きを生成する段階です

prefill は比較的大きな行列演算をまとめて流しやすく、GPU のスループットを出しやすいです。decode は 1 ステップごとの逐次依存が強いため、学習や prefill よりレイテンシの影響を受けやすくなります。

この比較表を頭に入れておくと、「推論が遅い」と言われたときに、どの段階が遅いのか を切り分けやすくなります。

推論サーバの内部イメージ

推論サーバが吸収しているもの

KV cache が重要な理由

KV cache は、過去トークンの attention 用キーとバリューを保存しておく仕組みです。これがあると、毎回すべての文脈を最初から計算し直さずに済みます。ただし代償として、

• 長いコンテキストほど VRAM を圧迫しやすい
• 小さなリクエストが多数来ると断片化しやすい
• マルチテナント環境では管理が難しい

という問題が出ます。vLLM の docs でも、PagedAttention、Automatic Prefix Caching、Quantized KV Cache、Speculative Decoding といった仕組みが前面に出ています。

つまり LLM 推論では、重みが主役なのはもちろんですが、運用に入ると KV cache 管理がもう一つの主役 になります。ここが画像分類や古典的な DNN 推論とかなり違うところです。

LLM 推論で GPU に効く資源

資源と症状の対応

なぜ「推論サーバ」が重要なのか

単体の GPU カーネルが速くても、実サービスではそれだけでは足りません。実際には、

• リクエストがばらばらのタイミングで到着する
• モデルごとに最適なバッチサイズが違う
• 会話の長さが偏る
• KV cache の寿命管理が必要
• 複数モデルや LoRA を共存させたい

といった運用課題があります。ここで、TensorRT、Triton Inference Server、vLLM のようなソフトウェアが効いてきます。

代表的な推論スタック

NVIDIA の最新 TensorRT docs では、TensorRT は NVIDIA GPU 上の deep learning inference を最適化・高速化する SDK とされ、FP32/FP16/BF16/FP8/INT8、dynamic shapes、transformer / LLM 最適化が明示されています。Triton docs では、HTTP または gRPC endpoint を通してモデル推論を提供するサーバと整理されています。vLLM docs では、PagedAttention、continuous batching、prefix caching、speculative decoding、OpenAI-compatible API server などが前面に出ています。

OpenAI 互換サービングの最小イメージ

from openai import OpenAI

client = OpenAI(
    base_url="http://localhost:8000/v1",
    api_key="dummy",
)

resp = client.chat.completions.create(
    model="my-llm",
    messages=[
        {"role": "user", "content": "GPU と CPU の違いを説明して"}
    ],
    max_tokens=128,
)

print(resp.choices[0].message.content)

このようなクライアント例が意味を持つのは、サーバ側で batching、KV cache、モデルロード、ストリーミング、スケジューリングを吸収してくれるからです。

量子化と GPU

量子化 は、重みや activations を低ビット化して、容量と帯域を削減する手法です。推論ではとくに重要です。

• FP16 / BF16: 比較的扱いやすい
• INT8: 容量削減と速度改善のバランスを取りやすい
• INT4 / FP4 / FP8 系: より攻めた圧縮だが、対応ハードウェアと品質検証が重要

量子化は「何 bit まで下げられるか」だけでなく、「その GPU とランタイムで速く動くか」が重要です。ビット数が低くても、実装が弱ければ逆に遅くなることがあります。

FlashAttention が効く理由

FlashAttention は、attention を単に数式として最適化するだけでなく、GPU のメモリ階層に合わせて再構成する代表例です。ポイントは、

• 大きな attention 行列をそのまま全部書き出さない
• タイル単位で計算して、共有メモリやレジスタで再利用する
• 帯域ボトルネックを緩和する

ことにあります。これは GPU 最適化全般にも通じる考え方で、「計算そのもの」より「中間結果をどこへ置くか」が勝負になる例です。

TensorRT の役割を分けて見る

TensorRT は、学習フレームワークそのものではなく、推論実行計画の最適化と実行 に強い層です。

• モデルグラフを解析する
• 演算を融合する
• 精度を下げられる場所を見つける
• GPU に合った実行 plan を作る

つまり TensorRT は、手書き CUDA の代わりではなく、「既存モデルを NVIDIA GPU 向けに本番用へ整える層」と見ると理解しやすいです。

vLLM が注目される理由

vLLM が強く注目されるのは、単に OpenAI 互換 API を出せるからではありません。実務では、

• PagedAttention による KV cache 管理
• continuous batching
• prefix caching
• speculative decoding
• 分散推論の整理

といった、「LLM サービングの痛点」へ直接答えていることが大きいです。GPU ハードウェアの性能差が小さい場面でも、こうしたランタイムの差で実効性能が大きく変わります。

speculative decoding

speculative decoding は、小さな補助モデルや予測機構で先読み候補を出し、大きな本体モデルでまとめて検証する発想です。うまくはまると decode の逐次性を少し崩せますが、

• モデルの組み合わせが難しい
• 外れが多いと得をしない
• メモリと実装が複雑になる

というトレードオフがあります。

推論運用で見るべきキュー

LLM 推論でのマルチ GPU

巨大モデルでは 1 枚に重みが載り切らないため、分割推論が必要になります。

• tensor parallel: 同じ層の計算を GPU 間で分ける
• pipeline parallel: 層の塊ごとに GPU を分ける
• data parallel: 独立リクエストを複数 GPU へ分散する
• context parallel などの派生手法: 長い入力を分散して扱う

推論では、学習ほど巨大な all-reduce が常に支配するわけではありませんが、境界の通信待ちは依然として大きな論点です。

LLM 推論でよくある失敗

• 重みは載るが KV cache で VRAM が足りなくなる
• 小さすぎるバッチで GPU が遊ぶ
• 長文リクエストが混ざって待ち行列が崩れる
• 量子化はしたが、実装が最適化されておらず速くならない
• サービス全体では CPU 前処理やトークナイザが詰まっている

LLM 推論を評価するときの観点

1. time to first token
2. tokens per second
3. 同時接続時のスループット
4. 長文混在時の tail latency
5. GPU メモリ効率
6. モデル切替や LoRA 混在時の挙動

ケーススタディ: 小さな LLM サービスが遅い

よくある状況として、「GPU を積んだのに応答が遅い」ケースがあります。たとえば、

• 同時接続は少ない
• 1 リクエストごとにすぐ実行している
• バッチングが弱い
• トークナイザと前処理が CPU に残っている

という状態では、GPU 側の演算性能が高くても全体は伸びません。この場合の改善順は、

1. time to first token と CPU 時間を分けて測る
2. continuous batching を導入する
3. KV cache の断片化と VRAM 使用率を見る
4. 量子化や prefix caching を検討する

となります。ここでは「GPU が遅い」のではなく、「GPU へ仕事を届ける流れ」が遅いことが多いです。

GPU 最適化の実践パターン

ここでは、実務で GPU カーネルや GPU ワークロードを改善するときの考え方を、手順ベースで整理します。GPU 最適化は魔法ではなく、計測 -> 仮説 -> 変更 -> 再計測 の繰り返しです。

全体最適から始める

最適化で最初にやるべきことは、最も遅い部分を正確に見つけることです。

実践パターン 1: 転送削減

もっともよく効くのに見落とされやすいのが、CPU と GPU の往復を減らすことです。

• 前処理を GPU 側へ寄せる
• 複数カーネルの間でデータを GPU 上へ保持する
• 小分け転送をまとめる
• pinned memory を使って転送効率を上げる

典型例

• 画像を CPU で decode -> normalize -> resize -> GPU 送信、ではなく一部を GPU 側へ寄せる
• 各ステップごとに CPU へ戻さず、後段のカーネルまで GPU メモリ上でつなぐ

悪い例と良い例

悪い例:
CPU preprocess -> GPU kernel A -> CPU copy back -> CPU preprocess -> GPU kernel B

良い例:
CPU preprocess once -> GPU kernel A -> GPU kernel B -> GPU kernel C -> CPU copy back once

実践パターン 2: アクセスの規則化

GPU は規則的なアクセスに強いので、データ配置を変えるだけで大きく伸びることがあります。

• Array of Structures を Structure of Arrays へ変える
• ランダムアクセスをソートやバケット分割で緩和する
• 近いスレッドが近いアドレスを触るよう並べ替える

AoS と SoA のイメージ

AoS:
[{x,y,z}, {x,y,z}, {x,y,z}]

SoA:
x = [x,x,x]
y = [y,y,y]
z = [z,z,z]

GPU では、多数スレッドが同じ属性を一斉に読むなら SoA の方が有利になりやすいです。

実践パターン 3: タイリング

タイリング は、グローバルメモリから読む大きなデータを小片へ分け、共有メモリやキャッシュで再利用しやすくする考え方です。行列積で頻出ですが、画像処理や stencil 計算でも重要です。

タイリングが効く理由

• グローバルメモリの再読み出しを減らせる
• 共有メモリの低レイテンシを活かせる
• 演算強度を上げやすい

タイリングの落とし穴

• 共有メモリを使いすぎて occupancy が落ちる
• bank conflict が増える
• タイル境界の処理が複雑になる

CUDA 風のタイリング断片

__shared__ float tileA[16][16];
__shared__ float tileB[16][16];

int row = blockIdx.y * 16 + threadIdx.y;
int col = blockIdx.x * 16 + threadIdx.x;
float acc = 0.0f;

for (int t = 0; t < K; t += 16) {
    tileA[threadIdx.y][threadIdx.x] = A[row * K + t + threadIdx.x];
    tileB[threadIdx.y][threadIdx.x] = B[(t + threadIdx.y) * N + col];
    __syncthreads();

    for (int k = 0; k < 16; ++k) {
        acc += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
    }
    __syncthreads();
}

C[row * N + col] = acc;

この例の本質は、「グローバルメモリから小片を読み、共有メモリへ置き、そこで再利用する」ことです。

実践パターン 4: カーネル融合

カーネル融合 は、連続する小さなカーネルをまとめて、途中結果をレジスタや共有メモリに留める考え方です。

• launch overhead を減らせる
• 中間結果のグローバルメモリ往復を減らせる
• ただしレジスタ pressure は増えやすい

融合のイメージ

実践パターン 5: 精度の見直し

混合精度や低精度化は、AI ワークロードで特に大きな武器です。

• FP32 から FP16 / BF16 へ下げる
• Tensor Core が効く形へ揃える
• INT8 や FP8 が使えるなら精度検証込みで試す

実践パターン 6: バッチ設計

GPU はある程度まとめて仕事を流した方が効率が出ます。バッチが小さすぎると、演算器も帯域も使い切れません。ただしレイテンシ要求が強い場合は、巨大バッチが正義ではありません。

推論バッチの擬似コード

pending = []

while True:
    req = maybe_pop_request()
    if req is not None:
        pending.append(req)

    if should_launch_batch(pending):
        batch = build_batch(pending)
        launch_gpu_inference(batch)
        pending = keep_unfinished_requests(pending)

実際のサービングランタイムはこれをはるかに複雑にしていますが、本質は「いつ束ねるか」と「どこで待たせるか」です。

実践パターン 7: 分岐の整理

ダイバージェンスが強いなら、

• 先に CPU 側で分類して近い要素をまとめる
• if を複数カーネルへ分ける
• predication や mask で揃える

といった方法が候補です。

実践パターン 8: ライブラリを疑う

手書きカーネルの前に、まず既存ライブラリが十分最適化されているかを確認します。

• GEMM は cuBLAS / rocBLAS
• DNN 基本演算は cuDNN / MIOpen
• 通信は NCCL / RCCL

これらで足りるなら、手書きは最後です。

最適化の失敗パターン

• occupancy だけ上げて満足する
• 単体カーネルだけ速くして end-to-end は変わらない
• ベンチマーク入力が実運用と違いすぎる
• コンパイラやドライバ更新で結果が揺れているのに気づかない
• GPU 利用率だけ見て「使い切れている」と誤解する

before / after で見る改善の思考

実務での進め方

1. 現状の end-to-end 指標を固定する
2. もっとも重い 1 箇所だけを選ぶ
3. 変更の狙いを 1 つに絞る
4. 効いたかどうかを数値で比較する
5. 次の 1 箇所へ進む

小さな最適化レポートの型

ベンチマーク結果の読み方

GPU ベンチマークは、単に ms が小さいかどうかだけでは不十分です。最低でも次の観点を一緒に見ます。

ベンチマークを比べるときの前提条件

比較には、少なくとも次を揃える必要があります。

• GPU 型番
• ドライバ / ランタイム / ライブラリ版
• バッチサイズ
• 精度形式
• 入力サイズ
• warmup 有無
• CPU 側前処理の含み方

このどれかが違うだけで、比較はかなり危うくなります。

ありがちな罠

• cold start を混ぜたまま比較する
• 異なる精度の結果を同列比較する
• 単発の最速値だけを載せる
• 実運用より小さすぎる入力で測る
• サービス系なのに p99 を見ない

ケーススタディ: 行列積カーネルが遅い

行列積でありがちな失敗は、「各要素ごとにその場で全部読む」実装です。これだと、

• グローバルメモリ読み出しが多すぎる
• 同じデータを何度も読み直す
• 演算器より先に帯域が詰まる

となります。改善の典型は、

1. タイリングして shared memory 再利用を増やす
2. block size を調整する
3. 可能なら Tensor Core が効く精度へ寄せる
4. 最後に cuBLAS などと比較する

です。ここでのゴールは「自作最適化が速いこと」ではなく、「既存ライブラリで十分か、自作が必要か」を判断することです。

ケーススタディ: 推論は速いがサービスは遅い

GPU 単体ベンチマークでは十分速いのに、実サービスの応答が悪いケースも多いです。この場合は、

• HTTP 層
• トークナイズ
• バッチ待ち
• モデル切替
• 出力整形

のどれかが支配していることがあります。ここで GPU カーネルだけを最適化しても、体感は変わりません。まず time to first token と GPU active time を分けるのが先です。

CUDA・ROCm・Metal・Vulkan のコードモデル比較

GPU プログラミングは API ごとに書き方が違いますが、根本の考え方はかなり共通です。ここでは「何が同じで、何が違うか」をコードモデルの観点から整理します。

4 つに共通する基本構造

どの API でも、概念的には次の 5 段階があります。

1. バッファや資源を確保する
2. GPU で走る関数を用意する
3. 実行単位の大きさを決める
4. コマンドをキューへ積む
5. 必要なら同期して結果を読む

少し大きいコード骨格で見る

最小ベクトル加算だけだと、実務で本当に悩む メモリ確保、資源束縛、dispatch 設定 の感覚が見えにくいです。そこで、実際に書くときの骨格を簡略化して並べます。

最小ベクトル加算で見る発想

CUDA / HIP 的な見え方

• __global__ なカーネルを定義する
• threadIdx, blockIdx, blockDim から自分の index を作る
• grid と block の大きさを起動時に決める

このモデルは「明示的にスレッド階層を考える」感覚が強いです。HIP もかなり近い書き味を取ります。

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, bytes);
cudaMalloc(&d_b, bytes);
cudaMalloc(&d_c, bytes);

cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

dim3 block(256);
dim3 grid((n + block.x - 1) / block.x);
vec_add<<<grid, block>>>(d_a, d_b, d_c, n);

cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

この書き味では、「自分でメモリ確保し、自分で起動形を決める」感覚が強いです。

Metal 的な見え方

• [[kernel]] 関数を定義する
• $thread_position_in_grid$ や threadgroup サイズを使う
• command buffer と compute encoder を作って dispatch する

Metal は Apple プラットフォーム全体との統合が強く、resource binding や pipeline state の扱いが API 設計に深く入っています。

#include <metal_stdlib>
using namespace metal;

kernel void vec_add(
    device const float* a [[buffer(0)]],
    device const float* b [[buffer(1)]],
    device float* c [[buffer(2)]],
    uint id [[thread_position_in_grid]]
) {
    c[id] = a[id] + b[id];
}

let commandBuffer = commandQueue.makeCommandBuffer()!
let encoder = commandBuffer.makeComputeCommandEncoder()!
encoder.setComputePipelineState(pipelineState)
encoder.setBuffer(bufferA, offset: 0, index: 0)
encoder.setBuffer(bufferB, offset: 0, index: 1)
encoder.setBuffer(bufferC, offset: 0, index: 2)
encoder.dispatchThreads(gridSize, threadsPerThreadgroup: tgSize)
encoder.endEncoding()
commandBuffer.commit()
commandBuffer.waitUntilCompleted()

Metal では、pipeline state と command buffer の存在感が強く、アプリ全体との統合感があります。

Vulkan Compute 的な見え方

• compute shader を用意する
• pipeline と descriptor set を作る
• command buffer に dispatch を積む

Vulkan は最も明示的です。柔軟ですが、学習コストは高めです。

#version 450
layout(local_size_x = 256) in;

layout(set = 0, binding = 0) readonly buffer A { float a[]; };
layout(set = 0, binding = 1) readonly buffer B { float b[]; };
layout(set = 0, binding = 2) writeonly buffer C { float c[]; };

void main() {
    uint i = gl_GlobalInvocationID.x;
    c[i] = a[i] + b[i];
}

Vulkan 側の骨格:
1. buffer を作る
2. descriptor set へ束縛する
3. compute pipeline を作る
4. command buffer に bind + dispatch を記録する
5. queue submit して完了を待つ

Vulkan は最も冗長に見えますが、そのぶん資源や同期を細かく制御できます。

4 方式を俯瞰する図

用語対応表

どこが一番つまずきやすいか

どう学ぶとよいか

1. まず CUDA か Metal のように学習導線が比較的明確なものを 1 つ学ぶ
2. スレッド階層、メモリ階層、同期の考え方を掴む
3. その後に HIP や Vulkan で概念対応を取る

どれを選ぶべきか

• NVIDIA 中心の AI / HPC: CUDA が第一候補
• AMD 含む移植や調達の自由度重視: HIP / ROCm
• Apple 向けアプリやローカル ML: Metal
• 長寿命なクロスプラットフォーム低レベル基盤: Vulkan Compute

コードを書く前に考えること

どの API を選んでも、本質は次の問いへ戻ります。

• 並列化できるか
• データはどう置くか
• 共有メモリや cache をどう活かすか
• 同期と分岐をどう減らすか
• 既存ライブラリで済むのか

API 選定の判断表

API 別の認知負荷

FAQ

GPU は CPU より何倍速い？

処理次第です。何十倍にもなることもありますが、ほとんど速くならないこともあります。GPU が速いのは、向いた計算に対してです。

VRAM が多ければ速い？

必ずしも速くはありません。容量は重要ですが、帯域、キャッシュ、実行器、ソフトウェア最適化も同じくらい重要です。

CUDA を学べば他にも応用できる？

できます。CUDA 固有の API はありますが、スレッド階層、メモリ階層、レイテンシ隠蔽、ダイバージェンスの感覚は他の GPU API にもかなり効きます。

GPU と NPU はどう違う？

GPU は汎用並列計算へかなり広く対応します。NPU はもっと特定の AI 演算へ特化することが多く、電力効率で強い一方、柔軟性では GPU が勝る場面が多いです。

まず何を学ぶとよい？

順番としては、

1. CPU と GPU の設計目的の違い
2. スレッド階層とメモリ階層
3. ダイバージェンスとコアレッシング
4. 小さなベクトル加算や行列積
5. プロファイリング

が素直です。

GPU プログラミングモデル実践ガイド

CUDA、ROCm、Metal、Vulkan Compute を実装の感覚で比較する

GPU プログラミングは API ごとに書き味が違います。ただし、根本の問いは共通です。どこへデータを置くか、どの単位でスレッドを束ねるか、どう dispatch するか、どこで同期するか。この実践ガイドでは、CUDA、HIP / ROCm、Metal、Vulkan Compute を、概念対応とコード骨格の両面から整理します。

目次

1. 共通して考えるべきこと
2. CUDA の見方
3. HIP / ROCm の見方
4. Metal の見方
5. Vulkan Compute の見方
6. 最小コード例
7. 少し大きいコード骨格
8. 用語対応表
9. どれを選ぶべきか
10. ケーススタディ
11. FAQ
12. 参考文献

共通して考えるべきこと

どの API でも、最終的に考えることは似ています。

1. データをどこへ置くか
2. GPU 上の関数をどう書くか
3. 実行単位をどう切るか
4. コマンドをどう投入するか
5. どこで同期するか

CUDA の見方

CUDA は、NVIDIA GPU 向けで最も普及した計算基盤です。

何が強いか

• 学習導線が比較的明確
• エコシステムが厚い
• cuBLAS、cuDNN、NCCL など周辺が強い

何に注意するか

• NVIDIA 依存
• 最初は動いても、最適化は別の難しさがある

HIP / ROCm の見方

HIP は CUDA に近い書き味の移植層で、ROCm は AMD GPU 向けの総合スタックです。

何が強いか

• CUDA 資産を AMD 側へ寄せやすい
• オープン寄りの選択肢を取りやすい

何に注意するか

• 「完全互換」と思い込みすぎない
• 実機とライブラリ差を必ず確認する

Metal の見方

Metal は Apple プラットフォーム向けの低レベル GPU API です。

何が強いか

• macOS / iOS / Apple Silicon と自然に統合される
• graphics と compute を一貫して扱いやすい

何に注意するか

• Apple 外では使えない
• feature family の理解が必要

Vulkan Compute の見方

Vulkan Compute は、Khronos 系の低レベルクロスベンダ API の compute 機能です。

何が強いか

• 非常に明示的で制御自由度が高い
• 長寿命なクロスプラットフォーム基盤になりやすい

何に注意するか

• 記述量が多い
• pipeline、descriptor、synchronization の明示性が重い

最小コード例

CUDA / HIP 風

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

Metal

#include <metal_stdlib>
using namespace metal;

kernel void vec_add(
    device const float* a [[buffer(0)]],
    device const float* b [[buffer(1)]],
    device float* c [[buffer(2)]],
    uint id [[thread_position_in_grid]]
) {
    c[id] = a[id] + b[id];
}

Vulkan Compute

#version 450
layout(local_size_x = 256) in;

layout(set = 0, binding = 0) readonly buffer A { float a[]; };
layout(set = 0, binding = 1) readonly buffer B { float b[]; };
layout(set = 0, binding = 2) writeonly buffer C { float c[]; };

void main() {
    uint i = gl_GlobalInvocationID.x;
    c[i] = a[i] + b[i];
}

少し大きいコード骨格

CUDA 的な骨格

float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, bytes);
cudaMalloc(&d_b, bytes);
cudaMalloc(&d_c, bytes);

cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

dim3 block(256);
dim3 grid((n + block.x - 1) / block.x);
vec_add<<<grid, block>>>(d_a, d_b, d_c, n);

cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

Metal 的な骨格

let commandBuffer = commandQueue.makeCommandBuffer()!
let encoder = commandBuffer.makeComputeCommandEncoder()!
encoder.setComputePipelineState(pipelineState)
encoder.setBuffer(bufferA, offset: 0, index: 0)
encoder.setBuffer(bufferB, offset: 0, index: 1)
encoder.setBuffer(bufferC, offset: 0, index: 2)
encoder.dispatchThreads(gridSize, threadsPerThreadgroup: tgSize)
encoder.endEncoding()
commandBuffer.commit()
commandBuffer.waitUntilCompleted()

Vulkan 的な骨格

1. buffer を作る
2. descriptor set へ束縛する
3. compute pipeline を作る
4. command buffer へ bind + dispatch を記録する
5. queue submit して完了を待つ

用語対応表

どれを選ぶべきか

ケーススタディ

ケース 1: 学習用コードを AMD へ持っていきたい

この場合は HIP / ROCm が候補になりますが、CUDA 完全互換前提で進めると危険です。依存ライブラリと性能差を早めに確認する必要があります。

ケース 2: Apple Silicon 上でローカル ML ツールを作る

Metal が自然です。Cross-platform 抽象化より、Apple 側の pipeline state や threadgroup 設計を理解する方が成果につながりやすいです。

ケース 3: エンジンレベルの compute 基盤がほしい

Vulkan Compute は有力ですが、学習コストはかなり高いです。長期保守を見据えた設計が必要です。

FAQ

まず何から学ぶのがよいですか

NVIDIA 環境なら CUDA、Apple 環境なら Metal が最も入りやすいです。

Vulkan は最強ですか

最も自由度は高いですが、最も書くことも多いです。常に最強ではありません。

HIP は CUDA の置き換えですか

かなり近いですが、完全に同一と考えるのは危険です。

GPU 最適化実践ガイド

memory bound、compute bound、転送、融合を end-to-end で判断する

GPU 最適化は、カーネルの中だけを見る作業ではありません。CPU との転送、バッチ設計、メモリ階層、同期、実行単位の設計、既存ライブラリの使い方まで含めて考える必要があります。この実践ガイドは、GPU を速くするというより、GPU ワークロードを正しく速くする ための実践整理です。

目次

1. 最適化の前提
2. GPU 最適化の手順
3. 転送最適化
4. メモリアクセス最適化
5. タイリングと再利用
6. カーネル融合
7. 分岐と同期の整理
8. 精度最適化
9. バッチ設計
10. ベンチマークの読み方
11. ケーススタディ
12. FAQ
13. 参考文献

最適化の前提

GPU 最適化は、次の誤解から外れるところから始まります。

• occupancy が高ければ勝ち、ではない
• GPU 使用率が高ければ勝ち、でもない
• カーネル単体が速ければサービス全体も速い、でもない

見るべき 3 層

1. end-to-end: ユーザーが待つ全体時間
2. runtime: CPU、転送、キュー、同期
3. kernel: GPU 上の純粋な計算

GPU 最適化の手順

基本原則

• 一度に 1 つずつ変える
• 変更前後を必ず同条件で測る
• 既存ライブラリで済むなら自作を急がない

転送最適化

CPU と GPU の往復は、非常に高価になりやすいです。

改善の方向

• CPU 前処理を GPU 側へ寄せる
• GPU 上で複数ステップを連鎖させる
• pinned memory を使う
• 小分け転送をまとめる

悪い:
CPU -> GPU -> CPU -> GPU -> CPU

良い:
CPU -> GPU -> GPU -> GPU -> CPU

メモリアクセス最適化

GPU は規則的なアクセスを好みます。

典型施策

• AoS から SoA へ変える
• 隣接スレッドが隣接アドレスを触るようにする
• ランダムアクセスを前処理で緩和する

AoS と SoA

AoS:
[{x,y,z}, {x,y,z}, {x,y,z}]

SoA:
x = [x,x,x]
y = [y,y,y]
z = [z,z,z]

タイリングと再利用

タイリング は、大きなデータを小片へ分け、shared memory や cache で再利用しやすくする手法です。

なぜ効くか

• グローバルメモリ再読込を減らす
• 演算強度を上げる
• 帯域律速を緩和する

タイリングの骨格

__shared__ float tileA[16][16];
__shared__ float tileB[16][16];

for (int t = 0; t < K; t += 16) {
    // global -> shared
    __syncthreads();
    // shared を再利用して演算
    __syncthreads();
}

落とし穴

• shared memory を使いすぎる
• bank conflict が増える
• 境界処理が複雑化する

カーネル融合

小さなカーネルを何本もつなぐと、launch overhead と中間結果の書き戻しが無視できなくなります。

何が嬉しいか

• launch overhead 減
• 中間結果のメモリ往復減

何が難しいか

• register pressure 増
• コードの複雑化

分岐と同期の整理

warp divergence や過剰同期は、GPU で非常に効率を落とします。

分岐で見ること

• 同じ warp 内で分岐が割れていないか
• 事前分類で揃えられないか
• 別カーネルへ分割した方がよくないか

同期で見ること

• block 内同期が多すぎないか
• 原子操作へホットスポットがないか
• CPU 側で GPU を細かく待っていないか

精度最適化

AI ワークロードでは、精度形式の変更が極めて強い最適化です。

バッチ設計

GPU はまとめて流す方が効率が出やすいですが、レイテンシ要件が強いと単純ではありません。

擬似コード

pending = []

while True:
    req = maybe_pop()
    if req:
        pending.append(req)

    if should_launch(pending):
        batch = build_batch(pending)
        run_gpu(batch)

見るべきこと

• 小さすぎて GPU が遊んでいないか
• 大きすぎて待ちが増えていないか
• 長文と短文が混ざって崩れていないか

ベンチマークの読み方

最低限見る指標

比較前提

• GPU 型番
• ドライバ / runtime
• 入力サイズ
• 精度形式
• batch size
• warmup

よくある誤読

• 最速 1 回だけ比較する
• 平均値だけで判断する
• CPU 前処理込みかどうかを混ぜる
• 精度が違うのに同列比較する

ケーススタディ

ケース 1: 行列積カーネルが遅い

疑う順は、

1. shared memory を使っているか
2. block size が極端でないか
3. Tensor Core が効く精度か
4. cuBLAS と比べてどれくらい差があるか

ケース 2: GPU 利用率は高いが遅い

利用率が高くても、memory stall だらけなら効率は悪いです。GPU 利用率だけではなく、帯域や stall 理由を見ます。

ケース 3: 単体カーネルだけ速い

サービス全体では CPU 前処理や転送が支配していることがあります。end-to-end の改善がなければ、最適化の価値は限定的です。

FAQ

occupancy は高いほどよいですか

高いほどよいとは限りません。手段であって目的ではありません。

手書き CUDA はいつ必要ですか

既存ライブラリで足りないとき、またはワークロードが特殊で十分な最適化が既存にないときです。

GPU 使用率だけ見れば十分ですか

十分ではありません。帯域、stall、転送、p95 / p99 を一緒に見る必要があります。

LLM 推論と GPU サービング実践ガイド

KV cache、連続バッチング、ランタイム最適化をつなげて理解する

LLM 推論は、GPU の価値が最も可視化される実務領域のひとつです。ただし実際の性能は、単純な FLOPS や VRAM 容量だけでは決まりません。prefill、decode、KV cache、continuous batching、量子化、GPU 間通信、推論サーバのスケジューリングまで含めて見てはじめて、なぜ速いのか、なぜ詰まるのかが見えてきます。

この実践ガイドで重視すること

• LLM 推論 = GPU カーネル ではなく、サービング全体の問題として理解する
• prefill と decode の違いを性能上の違いとしてつかむ
• KV cache、量子化、continuous batching がなぜ重要かを腹落ちさせる
• TensorRT、Triton、vLLM の役割分担を整理する
• time to first token と tokens per second を混同しない

目次

1. LLM 推論とは何か
2. GPU が LLM 推論で重要な理由
3. prefill と decode
4. KV cache
5. 推論サーバの内部構造
6. バッチングとスケジューリング
7. 量子化と精度戦略
8. FlashAttention と attention 最適化
9. TensorRT、Triton、vLLM の役割分担
10. マルチ GPU 推論
11. ベンチマークの見方
12. よくあるボトルネック
13. ケーススタディ
14. 実装・運用チェックリスト
15. FAQ
16. 参考文献

LLM 推論とは何か

LLM 推論は、学習済みの大規模言語モデルへ入力を与え、次トークンを順番に予測して応答を生成する処理です。ここでいう トークン は、単語そのものではなく、モデルが扱う離散的な単位です。

学習との違い

• 学習 は重みを更新する
• 推論 は重みを固定したまま出力だけを作る

推論は学習より軽く見えますが、実サービスでは

• 低レイテンシ
• 高スループット
• 同時接続
• 長文混在

を同時に満たしたくなるため、別の難しさがあります。

GPU が LLM 推論で重要な理由

LLM 推論の中核は、巨大な行列演算と attention です。これらは、

• 同種の演算を大量のデータへ適用する
• 高帯域メモリが効く
• 低精度演算ユニットを活かしやすい

ため、GPU との相性が極めてよいです。

GPU が効くが万能ではない理由

推論の全体は GPU だけで完結しません。

• トークナイズは CPU 側に残ることが多い
• API 層や認証、ルーティングは CPU 側
• バッチ待ちやキュー制御はサーバロジック側

つまり「GPU は速いのにサービスは遅い」という状況は普通に起きます。

prefill と decode

prefill は、最初に与えられた文脈をまとめて処理して内部状態を構築する段階です。decode は、その後 1 トークンずつ生成する段階です。

性能上の違い

• prefill は大きな GEMM をまとめて流しやすい
• decode は逐次依存が強く、帯域とレイテンシが効きやすい

このため、tokens per second が良くても time to first token が悪い、またはその逆、ということが起きます。

KV cache

KV cache は、過去トークンに対する attention のキーとバリューを保存する仕組みです。毎回すべてを再計算しないために不可欠です。

何が嬉しいか

• decode ごとの再計算を減らせる
• 長い会話を現実的な速度で扱える

何が難しいか

• 長文ほど VRAM を消費する
• リクエストが混在すると断片化しやすい
• マルチテナントでは誰の cache をどこまで残すかが難しい

PagedAttention 的な発想

vLLM docs で中心に置かれている PagedAttention は、KV cache を固定長ブロックへ分けて管理し、断片化と再利用効率を改善する考え方です。これは OS のページ管理に近い発想で、「連続巨大領域を毎回確保しない」ことが本質です。

推論サーバの内部構造

実際の推論サーバは、単なる model.generate() のラッパではありません。

典型的な構成要素

• リクエスト受付
• バッチ構築
• モデル選択
• GPU 実行
• KV cache 管理
• 出力ストリーミング

GPU だけ最適化しても、このどこかが弱いと全体は伸びません。

バッチングとスケジューリング

LLM サービングでは、いつバッチを組むか が極めて重要です。

なぜ難しいか

• バッチを大きくすると GPU は嬉しい
• でも待たせすぎるとレイテンシが悪化する
• 短文と長文が混ざると待ち行列が崩れる

continuous batching

continuous batching は、既存の decode ループへ新しいリクエストを継続的に混ぜる発想です。固定バッチより効率がよいことが多いですが、実装はかなり複雑になります。

pending = []
active = []

while True:
    req = maybe_get_request()
    if req:
        pending.append(req)

    active = refill_active_batch(active, pending)
    run_one_decode_step(active)
    active = drop_finished(active)

量子化と精度戦略

量子化は、重みや activation を低ビット化して容量と帯域を削減する手法です。

誤解しやすい点

• 低ビットなら必ず速いわけではない
• 品質が落ちないとは限らない
• GPU とランタイムの対応次第で効果が変わる

FlashAttention と attention 最適化

FlashAttention は、attention の計算順序を GPU のメモリ階層に合わせて組み替える代表例です。

なぜ効くのか

• 中間結果を丸ごとグローバルメモリへ書かない
• タイルごとに shared memory や register で再利用する
• 帯域ボトルネックを弱める

これは「数式を少し変える」のではなく、「GPU に合う実行形へ変える」最適化の好例です。

TensorRT、Triton、vLLM の役割分担

この 3 つはよく同列に語られますが、役割は違います。

ざっくりした見分け方

• TensorRT: GPU 上でどう速く走らせるか
• Triton: どうサービスとして出すか
• vLLM: LLM らしい痛点をどう吸収するか

マルチ GPU 推論

巨大モデルでは 1 枚に重みが載り切らないため、複数 GPU を使います。

主な分け方

• tensor parallel
• pipeline parallel
• data parallel
• context parallel

難しさ

• GPU 間通信
• トポロジ依存
• バランス崩れ
• デバッグの難しさ

ベンチマークの見方

ベンチマークは 1 回速かった だけでは不十分です。

比較時に揃える前提

• GPU 型番
• ドライバ
• ランタイム
• 精度形式
• バッチサイズ
• 入力長
• warmup 有無

よくあるボトルネック

• CPU トークナイズ
• バッチ待ち
• KV cache 断片化
• decode 側の帯域律速
• GPU 間通信
• モデルロードの揺れ

ケーススタディ

ケース 1: 小さな LLM サービスが遅い

• GPU は強い
• 同時接続は少ない
• でも応答が遅い

このときは GPU カーネルより、バッチ待ち、CPU 前処理、モデル warmup が原因のことが多いです。

ケース 2: VRAM は足りるのに接続数が伸びない

この場合は、断片化や KV cache 管理、長文リクエスト混在を疑います。

ケース 3: 単体ベンチは速いのに本番が遅い

本番では、

• API 層
• 観測
• 認証
• ストリーミング整形

が乗るため、単体カーネルの結果だけでは判断できません。

実装・運用チェックリスト

1. time to first token と tokens/s を分けて見る
2. KV cache の割当と断片化を監視する
3. バッチ戦略を固定せず、実トラフィックで見直す
4. 量子化は品質と速度を両方測る
5. マルチ GPU では通信も観測する
6. cold start と warm start を分けて評価する

FAQ

GPU が速いなら、なぜ LLM サービスは遅いのですか

GPU 以外の層が多いからです。CPU 前処理、バッチ待ち、KV cache、ネットワーク、サーバ実装が全部効きます。

VRAM が大きければ安心ですか

安心ではありません。容量だけでなく、帯域、断片化、KV cache の管理方法が同じくらい重要です。

vLLM を使えば全部解決しますか

しません。かなり強い選択肢ですが、モデル構成、GPU 構成、トラフィック特性との相性があります。

まとめ

GPU は、CPU とは異なる前提で設計された並列計算機です。スループット、メモリ帯域、実行単位、最適化、サービングまでを一続きで捉えると、向く処理と向かない処理の見分けがしやすくなります。