GPU

目次

主要項目のみを表示しています。詳細な小見出しは本文内で確認できます。

• 概要
• 先に押さえるCPUとGPUの違い
• GPUとは何か
• なぜGPUは速いのか
• GPUの歴史
• 基本アーキテクチャ
• グラフィックスパイプライン
• GPGPUとCUDA以後
• 実行モデル：スレッド、warp、wavefront
• メモリ階層
• 同期、ダイバージェンス、占有率
• テンソルコアと行列演算
• GPUプログラミングモデル
• 主要APIとエコシステム比較
• CPU・TPU・NPU・FPGAとの比較
• 主要ユースケース
• 性能最適化の勘所
• マルチGPUと相互接続
• 仮想化、分割、マルチテナンシー
• 運用・監視・トラブルシューティング
• GPUが向かない処理
• 現在の動向
• 判断の指針
• LLM推論とGPUアーキテクチャ
• GPU最適化の実践パターン
• CUDA・ROCm・Metal・Vulkanのコードモデル比較
• GPU基礎のFAQ
• GPUプログラミングモデル実践ガイド
• 共通して考えるべきこと
• CUDAの見方
• HIP / ROCmの見方
• Metalの見方
• Vulkan Computeの見方
• 最小コード例
• 少し大きいコード骨格
• 用語対応表
• どれを選ぶべきか
• GPU API選定のケーススタディ
• GPU API選定のFAQ
• GPU最適化実践ガイド
• 最適化の前提
• 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. 分岐や同期で実行束の足並みが崩れないか

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でまず速く作る方が全体最適になることがよくあります。

現在の動向

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計算」という意味ではとても重要です。

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

現在の実務では、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 のようなソフトウェアが効いてきます。

代表的な推論スタック

推論基盤は、単にモデルをGPUへ載せるだけではありません。TensorRTのような推論最適化、Tritonのような推論サーバ、vLLMのようなLLM向けランタイムを組み合わせ、FP32/FP16/BF16/FP8/INT8、dynamic shapes、batching、KV cache、prefix caching、speculative decodingなどを使って、速度、コスト、品質のバランスを取ります。

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別の認知負荷

GPU基礎の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 を、概念対応とコード骨格の両面から整理します。

共通して考えるべきこと

どの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して完了を待つ

用語対応表

どれを選ぶべきか

GPU API選定のケーススタディ

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

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

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

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

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

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

GPU API選定のFAQ

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

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

Vulkanは最強ですか

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

HIPはCUDAの置き換えですか

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

GPU最適化実践ガイド

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

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

最適化の前提

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前処理込みかどうかを混ぜる
• 精度が違うのに同列比較する

GPU性能最適化のケーススタディ

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

疑う順は、

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

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

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

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

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

GPU性能最適化の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 を混同しない

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間通信
• モデルロードの揺れ

LLM推論のケーススタディ

ケース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を分けて評価する

LLM推論のFAQ

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

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

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

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

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

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

まとめ

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

参考文献

公式・標準

• Pipelines and Shaders with Direct3D 12

講義・記事

• CUDA C++ Best Practices Guide
• CUDA C++ Programming Guide
• NVIDIA MIG Supported GPUs
• NVIDIA MIG User Guide
• ROCm Programming Guide
• Khronos Vulkan Guide
• What is SPIR-V
• What is Vulkan?
• Compute Shader Tutorial for Vulkan

解説・補助

• Compute Passes
• MTLComputeCommandEncoder
• NVIDIA GPU Operator with MIG
• NVIDIA TensorRT Documentation
• NVIDIA Triton Inference Server
• Apple Metal Resources
• ROCm Documentation
• vLLM Documentation