GPU(Graphics Processing Unit)のプログラムから性能を最大限に引き出すためのソフトウェアレベルの手法・技術群の総称。CUDA・OpenCL などの GPU プログラミング言語上でプログラマが手動で適用する変換・チューニングを主対象とし、コンパイラ自動最適化やハードウェア改修は含まない。 ## 定義 GPU は大量の並列スレッドを同時実行する throughput-oriented プロセッサであり、単一スレッドの性能より全体スループットを優先する。最適化とは、この並列性を活かしながらメモリ帯域幅・計算量・同期オーバーヘッドの制約を同時に満たすバランスを見つける行為である。 **GPU の基本実行モデル**: - スレッドは thread block に束ねられ、Streaming Multiprocessor (SM) にマップされる - SM 内スレッドは 32 スレッドの warp 単位で命令をロックステップ実行(SIMT) - メモリ階層: レジスタ(最速) → シェアードメモリ(SM 内) → L1/L2 キャッシュ → デバイスメモリ(低速) → ホストメモリ(PCIe 経由) - TLP(スレッドレベル並列性)と ILP(命令レベル並列性)の両方が利用率に寄与 ## 分類体系(Hijma ほか 2023 による) [[Pieter Hijma]] らは 450 本の論文を分析し、4 テーマ・28 技術に分類した([[@2023__CSUR__Optimization Techniques for GPU Programming]]): ### テーマ 1: メモリアクセス - **オンチップ**: 専用メモリ使用・ワープ関数・レジスタブロッキング・レジスタ使用削減・再計算 - **オフチップ**: [[コアレスドメモリアクセス]]・空間的ブロッキング・[[カーネルフュージョン]]・ソフトウェアプリフェッチング・データ圧縮・プリコンピュート ### テーマ 2: 不規則性 - ループアンローリング・[[分岐発散]]削減・疎行列フォーマット・カーネル分割・冗長計算削減 ### テーマ 3: バランシング - ベクトル化・高速数学関数・ワープ中心プログラミング・スレッド毎ワーク変化・スレッドブロックリサイズ・[[Auto-tuning]]・負荷分散・同期削減・アトミック削減・ブロック間同期 ### テーマ 4: ホストインタラクション - ホスト通信・CPU/GPU 計算分担 ## アプリケーション特性による最適化選択 カーネルの性格によって適用すべき最適化が異なる: - **compute-bound カーネル**: ループアンローリング・スレッド毎ワーク変化・ベクトル化・auto-tuning が有効 - **memory-bound カーネル**: コアレスドアクセス・空間的ブロッキング・レジスタブロッキング・カーネルフュージョンが有効 - **データ再利用あり**: 専用メモリ・空間的ブロッキング・レジスタブロッキング・カーネルフュージョンが特に有効 - **不規則カーネル**: データ圧縮・分岐発散削減・疎行列フォーマット・カーネル分割・負荷分散が有効 ## 横断的知見 - **Hopper の公式最適化方針は、カーネル内計算だけでなくデータ移動・同期・SM 間局所性を同時に扱う**: H100 の thread block cluster、DSMEM、TMA、非同期トランザクションバリアは、従来のコアレスドアクセスやカーネルフュージョンに加えて、複数 SM 間での転送と計算を重ね合わせる。H800 の実測が示すように、その効果は行列形状、ブロックサイズ、クラスタサイズに依存するため、機能の有無ではなく実行条件まで含めたチューニングが必要になる。(Source: [[@2022__NVIDIA Developer Blog__NVIDIA Hopper Architecture In-Depth]], [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **実運用の GPU 低利用率の主因はカーネルレベルの最適化不足ではなく、ホスト-GPU データ転送・バッチサイズ設定・チェックポイントの3つである**: [[@2024__ICSE__An Empirical Study on Low GPU Utilization of Deep Learning Jobs]] は Microsoft Platform-X の 400 実ジョブから 706 件の問題を発見し、最多はホスト-GPU データ転送の非効率(27.90%)、次いでバッチサイズ不適切(25.64%)、モデルチェックポイント(16.43%)と続く。Hijma ら([[@2023__CSUR__Optimization Techniques for GPU Programming]])が論じるコアレスドアクセス・カーネルフュージョン等のカーネル内最適化技術は未適用のまま、アプリケーション API の既定設定(DataLoader の `pin_memory=False`)やハイパーパラメータ選択(バッチサイズ)だけで大きな利用率低下が生じている。GPU プログラミング最適化の理論(何が速くなるか)と実運用の低利用率原因(何が遅くなっているか)は乖離している。(Source: [[@2024__ICSE__An Empirical Study on Low GPU Utilization of Deep Learning Jobs]], [[@2023__CSUR__Optimization Techniques for GPU Programming]]) - **最適化は相互依存する**: カーネルフュージョンはメモリアクセス削減だけでなくブロック間同期を可能にしてプリフェッチングを実現する。ループアンローリングはレジスタブロッキングを促進する。これらの連鎖があるため、単一最適化の効果を分離して評価することは困難である([[@2023__CSUR__Optimization Techniques for GPU Programming]]) - **採用頻度トップ 4 はコアレスドアクセス・専用メモリ使用・分岐発散削減・auto-tuning**: これらは GPU の基本的な制約(帯域幅・分岐・パラメータ探索)を直撃する技術であり、アーキテクチャが変化しても汎用性を保つ([[@2023__CSUR__Optimization Techniques for GPU Programming]] Fig. 3) - **アーキテクチャ世代によって効果が激変する**: Fermi の L1/L2 キャッシュ導入(2010)でコアレッシング効果が低下、Turing のキャッシュ統合(2018)でテクスチャメモリ効果が激減、Volta の Tensor Core(2017)で MMA 命令が特殊最適化ドメインを開いた([[@2023__CSUR__Optimization Techniques for GPU Programming]] §7.3) - **LLM 推論との接点**: [[LLM推論]]の高速化（[[FlashAttention]]・PagedAttention 等）はカーネルフュージョン・シェアードメモリ活用・Tensor Core 利用といった GPU 最適化技術の直接応用である。ただし本サーベイ（2021 年以前）は LLM 時代の最適化を対象外としている - **FlashAttention シリーズは「ボトルネック追跡型最適化」の教科書的事例**: [[FlashAttention]] の 4 世代にわたる進化は、Hijma ら([[@2023__CSUR__Optimization Techniques for GPU Programming]])の 4 テーマ分類体系における「メモリアクセス」テーマ（カーネルフュージョン・専用メモリ使用・空間的ブロッキング）と「バランシング」テーマ（ワーク分割・同期削減）の技術が、GPU アーキテクチャの世代交代に伴ってどのように再組み合わせされるかの実例を提供する。FA1 はテーマ 1 のカーネルフュージョンが主、FA2 はテーマ 3 のワーク分割が主、FA3-4 はテーマ 1+3 の非同期パイプラインが主であり、最適化の「効く場所」が世代ごとに移動している。(Source: [[@2022__arXiv__FlashAttention - Fast and Memory-Efficient Exact Attention with IO-Awareness]], [[@2026__arXiv__FlashAttention-4 - Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling]], [[@2023__CSUR__Optimization Techniques for GPU Programming]]) - **LLM 推論のボトルネック診断は Hijma らの分類が想定する静的な「最適化技術の選択」から、Nsight Compute/Systems による動的診断へ重心を移す**: LLM高速化の勉強会資料は、CUDA カーネル単体の性能検証(メモリストール vs SM 飽和)に Nsight Compute を、システム全体のボトルネック特定に Nsight Systems + NVTX を使うと説明し、CUDAGraph 未適用時のカーネル間 100 μs ギャップという具体的な観測例を示す。これは「実運用の GPU 低利用率はカーネル最適化不足ではなくホスト-GPU 転送やバッチサイズ設定に起因する」という Microsoft の実証研究と同じ問題意識——理論的な最適化技術のカタログだけでは実際のボトルネック箇所を特定できず、プロファイラによる実測が不可欠である——を LLM 推論の文脈で裏付ける。(Source: [[@2024__ICSE__An Empirical Study on Low GPU Utilization of Deep Learning Jobs]], [[@2026__SpeakerDeck__LLM高速化(勉強会)]]) - **GPU 利用率低下の原因としてのホスト-GPU データ転送は、ストレージ I/O データパスの選択次第でさらに増減しうる**: [[@2024__ICSE__An Empirical Study on Low GPU Utilization of Deep Learning Jobs]] はホスト-GPU データ転送の非効率を GPU 低利用率の最多原因(27.90%)とするが、そのデータ転送の入口であるストレージ I/O パス自体の選択も性能を左右する。[[@2026__POMACS__Towards Scalable Storage Architectures for GPU Clusters Running Large Language Models]] は、事前学習・ファインチューニングの大規模シーケンシャル転送では [[GPUDirect Storage]] が CPU 使用率を libaio 比で約 1/4、interrupt-driven io_uring 比で約 1/8 に削減しつつ同じ帯域上限に到達することを示す。GPU 利用率の理論(カーネルレベルの計算最適化)とは異なる層——ストレージからホスト DRAM・GPU メモリへの経路そのもの——にも独立した最適化余地があることを示唆する。(Source: [[@2024__ICSE__An Empirical Study on Low GPU Utilization of Deep Learning Jobs]], [[@2026__POMACS__Towards Scalable Storage Architectures for GPU Clusters Running Large Language Models]]) - **LLM エージェントは本分類体系の一部の技術にしか到達できていない**: [[@2026__arXiv__FlashInfer-Bench - Building the Virtuous Cycle for AI-driven LLM Systems]] の評価では、LLM エージェントは高レベル DSL(Triton)ではコンパイラ任せで auto-tuning・命令選択(テーマ3・Hijma らの分類)の恩恵を受けられるが、CUDA では専用メモリ活用・レジスタブロッキングなど手動最適化(テーマ1)を要する場面でハードウェア仕様を提示されても活用に失敗する例が多い。人間の GPU 最適化技術カタログ(本ページ)のうちどの技術がエージェントに「自動的に」届き、どの技術が依然として明示的な人間の介入を要するかを対比する軸として [[LLM駆動GPUカーネル生成]] を参照。(Source: [[@2026__arXiv__FlashInfer-Bench - Building the Virtuous Cycle for AI-driven LLM Systems]]) - **GPU 最適化の対象はカーネル内部から、ホスト、通信、メモリ、サービング設定を含むワークロード全体へ広がる**: [[Zoomer]] は GPU 利用率、メモリ帯域、通信パターン、CPU／GPU 律速、rank 間負荷不均衡、推論の要求レイテンシを同時に分析する。既存のカーネル最適化技法が「どの実装変換が効くか」を扱うのに対し、Zoomer はまず実行トレースからボトルネック層を特定し、一行のメモリコピー修正やサービングパラメータ調整までを改善候補として提示する。(Source: [[@2025__EngineeringAtMeta__Zoomer - Powering AI Performance at Meta's Scale Through Intelligent Debugging and Optimization]]) - **アーキテクチャ固有の最適化は、ワークロードの計算密度に合わせて選ぶ必要がある**: Hopper の評価では、`wgmma` は `N >= 64` の大きな行列でピークに近づく一方、`N < 64` では共有メモリ待ちを隠せず性能が低下した。非同期コピーも小さいブロックでは改善するが、大きいブロックでは同期コピーが待ち時間を隠して優位になる場合がある。(Source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **アーキテクチャ最適化の効果は命令レベルから実ワークロードまで一貫しない**: Blackwell の FP4/FP6 MMA は GB203 の低精度・高 ILP に適する一方、同じ GB203 の大規模 FP8 D-GEMM は H100 の約4分の1のスループットにとどまった。最適化対象は Tensor Core の対応形式だけでなく、タイル形状、カーネル選択、メモリ階層、ワープスケジューリング、電力状態まで含める必要がある。(Source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]], [[@2025__arXiv__Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks]]) - **低精度化は性能・メモリ・電力を同時に改善するとは限らない**: Blackwell の FP4 は Tensor Core マイクロベンチマークで16.753 Wと低いが、FP8 D-GEMM では Blackwell が低スループットと最大114.4 Wの電力ピークを示した。低精度最適化は形式単体でなく、実ワークロードと性能毎ワットで検証する必要がある。(Source: [[@2025__arXiv__Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks]]) - **GPU 最適化はカーネル内の演算変換だけでなく、データ表現とホスト間経路の設計から始まる**: 2004 年の GPU クラスタ実装は、LBM の速度分布を 2D テクスチャへ詰め、分散した境界データを 1 テクスチャへ集約してから GPU–CPU 転送することで読み出し回数を抑えた。現代の GPU 最適化が共有メモリ・非同期コピー・カーネル融合を重視するのに対し、同じ「データ移動を計算の前に設計する」という原則がテクスチャメモリと AGP の時代にも現れている。(Source: [[@2004__SC__GPU Cluster for High Performance Computing]], [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **通信と計算のオーバーラップは、クラスタのスケール限界を消すのではなく限界点を後ろへ移す**: Fan ほかは LBM で 28 ノードまではネットワーク通信を計算へ重ね合わせたが、28 ノード以上では非重複通信が増え、32 ノードの効率は 66.8% になった。Hopper の単一 GPU 内非同期機構と大規模 LLM 訓練の通信オーバーラップを並べると、隠蔽可能な待ち時間と残余クリティカルパスを分けて測る必要がある。(Source: [[@2004__SC__GPU Cluster for High Performance Computing]], [[@2024__NSDI__MegaScale - Scaling Large Language Model Training to More Than 10,000 GPUs]], [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **アーキテクチャ世代が変わると最適化対象の層は変わるが、通信量と計算量の比を管理する原理は残る**: LBM はサブドメインを立方体に近づけて境界面積/体積比を下げ、GPU クラスタの通信量を抑えた。後年の GPU クラスタでは並列化方式・ラック内外トポロジ・集団通信を協調設計するが、計算を細かく分割するほど境界・通信の割合が増えるという制約は共通する。(Source: [[@2004__SC__GPU Cluster for High Performance Computing]], [[@2024__NSDI__MegaScale - Scaling Large Language Model Training to More Than 10,000 GPUs]], [[@2026__MLSys2026__SAKURAONE - An Open Ethernet-Based AI HPC System]]) - **共有メモリ型の統合グラフィックスでも、GPU最適化の中心はワークロード別の資源均衡である。** - Intel 915の設計は、インタポレータ数、L1テクスチャ・キャッシュのポート数、Pixel Shader FPU数、Pixel Processingの並列度をベンチマークとシミュレーションで選択する。（Source: [[@2005__Intel Technology Journal__High-Performance Graphics and TV Output Comes to the Second-Generation Intel Centrino Mobile Technology Platform]]） - **GPU最適化では、電力・熱制約を性能指標と同じ設計変数として扱う必要がある。** - 第2世代モバイルプラットフォームの統合グラフィックスは、性能向上とバッテリー寿命・熱管理を同時に成立させる対象として説明される。（Source: [[@2005__Intel Technology Journal__Second-Generation Intel Centrino Mobile Technology Platform]], [[@2005__Intel Technology Journal__Interface Material Selection and a Thermal Management Technique in Second-Generation Platforms Built on Intel Centrino Mobile Technology]]） ## 未解決の問い - Ampere(2020)・Hopper(2022)以降のアーキテクチャ(HBM3・NVLink・Transformer Engine)において各最適化技術の効果はどう変化するか - LLM 推論に特化した最適化技術（[[FlashAttention]]・PagedAttention・continuous batching）は本分類体系のどのテーマに属し、どう拡張されるか。FA1-4 の進化から、テーマ 1（メモリアクセス）とテーマ 3（バランシング）の組み合わせが主であることは見えたが、テーマ 2（不規則性）との交差点はあるか - [[Auto-tuning]] フレームワークは探索空間の爆発をどう扱うか。機械学習ベースのチューナーは汎化するか - 複数最適化を同時に適用したときの相互作用を事前に予測するモデルは存在するか - ジョブ投入前にバッチサイズ・チェックポイント頻度・DataLoader 設定の組み合わせで得られる GPU 利用率を予測するツールは、実運用でどれほど有効か。静的解析(コードアドバイザー)と動的プロファイリングのどちらが費用対効果が高いか([[@2024__ICSE__An Empirical Study on Low GPU Utilization of Deep Learning Jobs]]が提起する未解決方向) - Hopper 以降の GPU で、行列形状・ブロック形状・クラスタ形状を入力として `wgmma`、TMA、分散共有メモリの選択を自動化するチューナーをどう設計するか。 - Blackwell の FP4/FP6 と `mma`/`tcgen05` の選択を、行列形状・ILP・ワープ数・電力制約から自動化するチューナーは設計できるか。 - 規則格子アプリケーションのデータ配置・サブドメイン形状・GPU–CPU 転送・ノード間通信を、現代の HBM/NVLink/RDMA の階層へ自動写像する最適化器をどう設計するか。