# GPU占有率(Occupancy) ## 定義 GPU占有率(occupancy)は、ストリーミングマルチプロセッサ(SM)が同時に保持できるアクティブワープ数の理論上限に対する、実際にアクティブなワープ数の割合である。GPUはSIMT(Single Instruction, Multiple Threads)実行モデルの下で多数のワープ(32スレッド単位)を切り替えながら実行することでメモリアクセスや命令のレイテンシを隠蔽する。占有率が高いほど、あるワープがメモリロードやキャッシュフィルで停止しても別のワープに切り替えて実行を継続できるため、SMの計算ユニットを継続的に稼働させやすい。ただし占有率はレジスタ数・共有メモリ使用量などスレッドあたりの資源制約とのトレードオフにあり、無制限に上げられるわけではない。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] §Threads, Warps, Blocks, and Grids) ## 占有率を左右する要因 - **スレッドブロックサイズ**: ワープサイズ(32)の倍数でないと、部分的にしか埋まらないワープがスケジューラ枠を無駄に占有する。Blackwellでは256〜512スレッド/ブロックが占有率とリソース制約のバランスとして推奨される。 - **SMごとのハードウェア上限**: Blackwellでは1SMあたり最大64常駐ワープ(2,048スレッド)・最大32常駐ブロックという上限があり、これを超える要求はコンパイラ/ランタイムに縮小される。 - **レジスタ・共有メモリ使用量**: スレッドあたりのレジスタ使用量が増えるほど、1SMに同時常駐できるスレッド数が減り占有率が下がる。上限を超えるとレジスタスピル(ローカルメモリへの退避)が発生し、さらに性能を落とす。 - **`__launch_bounds__`とCUDA Occupancy API**: `__launch_bounds__(maxThreadsPerBlock, minBlocksPerSM)`アノテーションやランタイムの`cudaOccupancyMaxPotentialBlockSize()`により、コンパイラのレジスタ割当・インライン化判断を占有率志向に誘導できる。 ## 横断的知見 - **占有率の最大化と性能最大化は同義ではない**: 書籍(ch.6)は、逐次実装から並列実装への変更で占有率1.3%→38.7%・実行時間22倍短縮という劇的な改善例を示す一方、章末の結論では「占有率が高くても命令レベル並列性(ILP)が十分ならより低い占有率でも高スループットを達成でき、逆にアクティブスレッド数を減らしてレジスタを解放した方がスループットが上がる場合もある」と明確に釘を刺している。占有率は診断指標として有用だが、最終目標はスループットであり、占有率はその代理指標に過ぎない。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] §Maintaining High Occupancy and GPU Utilization, §Conclusion) - **占有率向上とRoofline上のメモリバウンド脱却は別軸の最適化である**: 占有率はレイテンシ隠蔽(=同一メモリバウンド状態でもGPUをアイドルにしない)の手段であるのに対し、[[Rooflineモデル]]が示す演算強度の引き上げ(低精度化等)はメモリバウンドという状態そのものを脱却する手段である。両者は独立に効き、書籍のLLM decodeフェーズの議論(メモリバウンドな重み転送)は、占有率を上げても解決しない典型例として位置づけられる。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] §Maintaining High Occupancy and GPU Utilization, §Roofline Model) ## 未解決の問い - `__launch_bounds__`やOccupancy APIが提案する設定は、実際のカーネルで±1〜2候補のブロックサイズを実測して検証すべきとされるが、この検証をCIやオートチューニングパイプラインに組み込む標準的な方法論は確立されているか。 - 占有率とレジスタスピルのトレードオフの「最適点」は、カーネルの種類(メモリバウンド/演算バウンド)によってどの程度体系的に予測できるか。それとも実測に頼らざるを得ないか。 - LLM推論のdecodeフェーズのようにメモリバウンドが本質的なワークロードにおいて、占有率チューニングはどこまで有効で、どこから低精度化・カーネルフュージョンなど他の最適化に切り替えるべきかの判断基準はあるか。 ## 関連 - 概念: [[CUDA]] / [[GPU最適化]] / [[メモリ階層とキャッシュ]] / [[Rooflineモデル]] - ソース: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] ## 出典 - [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]](§Threads, Warps, Blocks, and Grids / §Choosing Threads-per-Block and Blocks-per-Grid Sizes / §Maintaining High Occupancy and GPU Utilization / §Tuning Occupancy with Launch Bounds / §Conclusion)