2026-08-06·無料·12分で読めます

GPUはなぜ速いのか — 実行モデルと並列化の限界

CPUとGPUは「速い」の定義が違います。トランジスタ予算の使い道、32スレッドを束ねるSIMTとワープ、分岐発散で性能が落ちる理由、占有率とレジスタ圧。最後はAmdahlの法則で「どこまで速くなるか」を着手前に見積もり、プロファイラのどの指標を見てCPU律速に気づくかまで降ります。

要点

CPUはレイテンシ最適化、GPUはスループット最適化で、トランジスタ予算を制御に使うか演算器に使うかが違う。
GPUはレイテンシを縮めず、数万スレッドを常駐させて覆い隠す。これはリトルの法則の N そのものである。
SIMTは連続する32スレッドの束に同じ命令を流す方式で、分岐が割れると時間は通った分岐先の和になり、最悪32倍になる。
占有率はレジスタ・共有メモリ・ハードの固定上限で決まるが、目的ではなく症状の1つとして読む。
Amdahlの法則 S(p)=1/((1-f)+f/p) の上限は 1/(1-f) で、5%が直列なら無限に並べても20倍で頭打ちになる。

#gpu #simt #warp #occupancy #amdahl #parallel

対象textタスクnumerical

比喩: 名人ひとりか、平凡な百人か

同じ店を2通りに作れます。ひとつは、動線を完璧に頭に入れた名人をひとり置く店。次に何が要るかを先読みし、手を止める時間を極限まで削ります。もうひとつは、平凡な料理人を百人並べる店。ひとりひとりは遅く、誰かは必ずオーブンの前で突っ立っている。ところが百人いれば、誰かが待つあいだに別の誰かが動けます。待ち時間は一切消えていないのに、店全体が出す皿の数では後者が勝つ。

CPUとGPUの違いはこれです。CPUは1つの仕事を短い時間で終わらせる機械——レイテンシ最適化。GPUは単位時間あたりの仕事量を最大化する機械——スループット最適化。「速い」の定義が違うので、片方の物差しでもう片方を測ると必ず誤解します。

トランジスタ予算の使い道

チップに載るトランジスタは有限で、設計とはその予算をどこへ配るかという意思決定です。

CPUは大半を、1本の命令列を止めずに流すことへ使います。分岐予測器、投機実行、リオーダバッファ、大きなキャッシュ——これらは1つも演算をしません。すべて「1つのスレッドが待たされないように」するための投資です。

GPUは逆張りをしました。予測も投機もほとんどせず、その面積を演算器とレジスタファイルに回す。GPUの演算単位（NVIDIAならSM、ストリーミング・マルチプロセッサ）のレジスタファイルは数百KB規模で、同じチップのL1キャッシュより大きいことも珍しくありません。容量が欲しいのではなく、大量のスレッドの状態を退避も復帰もせずに常駐させるためです。おかげで切り替えはほぼ無料になり、「メモリ待ちが起きたら、待つあいだ別のスレッドを走らせる」が現実的な戦略になります。

メモリの壁のリトルの法則を思い出してください。帯域を埋めるには $N = \lambda L$ 件のアクセスを常に空中へ浮かせる必要がありました。GPUの数万スレッドは、その $N$ そのものです。 キャッシュでレイテンシを縮めるのではなく、並列度で覆い隠す。ここがGPU設計の芯です。

先に釘: 並列化が買えるのは定数倍だけ

ただし並列化で買えるのは定数倍です。 $p$ 台使って $p$ 倍が上限で、計算量のオーダーは1段も下がりません。 $O(n^2)$ は千台あっても $O(n^2)$ のまま、 $n$ を10倍にすれば百倍になって返ってきます。

FIG 1並列化は曲線を下へ平行移動させるだけで、形は変えない。千台並べても指数の曲線にはすぐ追い抜かれる——「まずアルゴリズム、次に並列化」の順序が変わらない理由

順序は動きません。まずオーダーを下げ、次にデータの動かし方（行列積のコスト）、それから並列化。後半のAmdahlの法則は、その最後の一手が何倍で頭打ちになるかを着手前に教えてくれる道具です。

SIMT — スレッドに見えて、動くのは束

大量のスレッドを本当に独立に走らせると、制御回路の節約になりません。命令を取って解読する回路が、スレッドの数だけ要るからです。

そこでGPUは妥協します。連続する32スレッドを1つの束にまとめ、束の全員へ同じ命令を流す。 NVIDIAではこの束をワープ（warp）、AMDではウェーブフロント（wavefront）と呼びます（幅は世代により32または64）。フェッチとデコードは束につき1回、演算器は32レーンぶん並び、同じ命令をそれぞれ別のデータに当てます。

これが SIMT（Single Instruction, Multiple Threads）です。プログラマにはスレッド単位の普通のコードを書かせながら、ハードは幅32のSIMDとして動く。書きやすさと効率のあいだに置かれた妥協であり、落とし穴の多くはこの継ぎ目から噴き出します。代表が、「束の全員に同じ命令」という前提が崩れる瞬間です。

分岐発散 — 束の中で意見が割れたとき

ワープの32レーンが if (x > 0) に差しかかり、17レーンが真、15レーンが偽だったとします。ハードは1つの命令しか流せないので、両方は同時に走りません。まず真の側を流し、偽だった15レーンはマスクで無効化して結果を捨てる。次に偽の側を流し、今度は17レーンぶんを捨てる。

T_{\text{warp}} = \sum_{p \,\in\, \text{通った分岐先}} T_p

(1)

「ワープの実行時間は、通った分岐先の実行時間の和」です。つまり、17人が右、15人が左へ行きたい交差点に来たバスのようなもの。バスは右へ走って戻り、それから左へ走ります。 $T_p$ は分岐先 $p$ を1回通り抜けるのにかかる時間で、その道を選んだのが何人であっても、通った以上は32人全員がその時間ぶん座って待つ、ということです。1本で済んだところが $T_{\text{if}} + T_{\text{else}}$ になり、どちらの区間でも演算器の一部は捨てる結果を作るために動いています。この無駄を測るのがワープ実行効率 $E$ ——1命令あたり有効だったレーン数の平均を32で割った値で、発散がなければ $E = 1$ 、半分に割れれば $E = 0.5$ です。

最悪ケースはレーンごとに違う枝へ散る場合。スレッド番号で32方向へ分岐すれば、32本の枝を順に実行してそれぞれ1レーンしか有効になりません。効率 $1/32$ 、時間は32倍。条件の書き方だけで32倍損をすることがあり得ます。

ループも同じで、反復回数がスレッドごとに違えば最も長いスレッドが終わるまで束全体が拘束されます。可変長の系列を1スレッド1系列で処理すると、長いものが1本混じるだけで残り31本が引きずられる。長さでソートしてから束ねる定石はここから来ています。

なおVolta世代以降のNVIDIA GPUは各スレッドが独立したプログラムカウンタを持ち、発散したまま前進できます。細かい同期を安全に書けるようにした改良で、同時に走るのは同じ命令のレーンだけという事実は変わりません。

占有率とレジスタ圧

在庫の話に戻ります。演算単位に同時常駐できるワープ数が在庫で、その充足率が占有率（occupancy）。上限はレジスタ・共有メモリ・ハードの固定上限で決まります。レジスタから見ると、演算単位のレジスタ総数を $F$ 、1スレッドが使うレジスタ数を $R$ として

W \le \left\lfloor \frac{F}{32R} \right\rfloor

(2)

$W$ は常駐できるワープ数、32はワープの幅です。つまり、レジスタは演算単位ごとに数の決まった共有の棚で、1スレッドがそこから $R$ 個ずつ持ち出していく。棚にある $F$ 個を、ワープ1本ぶんの取り分（32スレッド× $R$ 個）で割って、丸ごと1本ぶん確保できた回数だけワープが居座れる、ということです。端数は切り捨て——1個足りなければ、その1本は入れません。 $R$ が増えると $W$ は階段状に落ちます。ローカル変数を1つ足しただけで常駐ワープが半分になり性能が崖から落ちる——これがレジスタ圧。さらに悪いのがスピルで、溢れた変数は「ローカルメモリ」という名の領域へ退避されますが、実体はDRAMです。

ただし占有率は高いほど良い、ではありません。目的はレイテンシを隠すことであって、ワープ数を増やすことではない。1スレッドが独立した複数のロードを同時に発行できるなら（命令レベル並列性）、少ないワープでも在庫は足ります。占有率のためにレジスタを削って逆に遅くなるのはよくある話で、占有率は原因ではなく症状の1つとして読むのが正確です。

Amdahlの法則 — 着手前に上限を出す

プログラムのうち割合 $f$ が完全に並列化でき、残り $1-f$ はどうしても直列だとします。 $p$ 並列にしたときの速度向上は

S(p) = \frac{1}{(1-f) + \dfrac{f}{p}}

(3)

$S$ は実行時間を何分の1にできるかの倍率、 $f$ は並列化できる部分の時間比（コード行数ではありません）、 $p$ は並列度です。つまり、仕事のうち手分けできる部分だけが人数で割られ、残りは何人集めようと誰かひとりが最初から最後まで自分でやるしかない、ということ。分母に居座る $(1-f)$ が、その「どうしても縮まない床」です。 $p$ をいくら増やしても $f/p$ が0へ近づくだけなので、上限は $1/(1-f)$ 。5%が直列なら、演算器を無限に並べても20倍で頭打ちです。

$f = 0.95$ 、 $p = 100$ なら $S \approx 16.8$ 倍。83台ぶんは捨てているわけです。使い道は2つ。着手前の足切りとして対象部分の時間比を先に測ること。そして逆に解いて $f$ を推定すること——2つの並列度で実測すれば $f$ が求まり、あとどこまで行けるかが見えます。

ただしAmdahlは問題サイズを固定したときの話です。台数を増やしたぶん問題を大きくできるなら（Gustafsonの見方）上限はここまで厳しくなく、学習が台数に素直にスケールするのはたいていこちら側だからです。

現場ではこう使う

プロファイラで最初に見る3つ

止まっている理由（Warp State Statistics）。Nsight Compute（ncu --set full）が、どの命令で何サイクル待たされたかを理由別に出します。Stall Long Scoreboard はグローバルメモリ待ち、Stall Barrier は __syncthreads() 待ち＝ブロック内の負荷不均衡、Stall MIO Throttle は共有メモリや特殊関数ユニットの詰まり。どれで止まっているか分かるまでコードを触らない。
ワープ実行効率（smsp__thread_inst_executed_per_inst_executed）。32に近ければ発散なし、16なら半分捨てています。分岐発散を疑うときの直接証拠です。
実測占有率（sm__warps_active.avg.pct_of_peak_sustained_active）。理論値ではなく実測を見る。乖離していたらブロック数不足かテール効果です。

コンパイル時にも情報が出ます。nvcc -Xptxas -v はカーネルごとのレジスタ数とスピルバイト数を印字します。スピルが0でないなら、他を見る前にそこから。

ブロックサイズの決め方

ワープ幅の倍数にして、まず128か256から。32の倍数でないと最後のワープに遊びレーンが出ますし、小さすぎるとブロックごとの固定費（起動、同期）が重く、大きすぎると1ブロックが資源を占有して常駐ブロック数が減ります。
迷ったら cudaOccupancyMaxPotentialBlockSize に計算させる。実際のレジスタ数と共有メモリ量から、占有率が最大になる値を返します。
占有率を上げたいときは __launch_bounds__(256) を付ける。レジスタ割り当てを抑えさせる宣言ですが、抑えすぎるとスピルするので -Xptxas -v で確認します。
グリッドは演算単位の数の整数倍を意識する。ブロック数が1.1倍のような半端だと、最後の1割のあいだ大半の演算単位が空になります（テール効果）。

CPU律速に気づく方法

nvidia-smi のGPU利用率は効率ではありません。「その瞬間カーネルが1本でも載っていたか」を集計するだけなので、極端に非効率なカーネルでも100%と出ます。見るべきはタイムライン。Nsight Systems や torch.profiler でカーネルの帯を並べ、カーネルとカーネルのあいだの空白を探します。空白の時間は、ホスト側にいます。

原因はたいてい3つ。①データローダが間に合っていない（num_workers、pin_memory、前処理のGPU移送を疑う）。②暗黙同期。.item()、.cpu()、print(tensor)、torch.cuda.synchronize() はそこでパイプラインを止めます。ループ内のロギング1行で非同期実行の利点が消えることがある。③カーネルが細かすぎる。起動には毎回数マイクロ秒の固定費があるので、数十マイクロ秒のカーネルを何千本も並べると起動費が支配します。融合（torch.compile）やCUDA Graphsで起動回数を減らします。

面接・設計レビューで出る問い

「並列化率95%のコードを100台で回すと何倍ですか」——16.8倍、上限は20倍。「ワープ内でスレッド番号の偶奇で分岐したら？」——枝を順に実行するので時間は約2倍、効率は50%。「占有率を100%にすれば最速ですか」——いいえ、レイテンシが隠せていれば十分です。どれも答えの数字より、律速している資源を言い当てられるかを見られています。

まとめ

CPUはレイテンシ最適化、GPUはスループット最適化。予算を制御に使うか演算器に使うかの違い
GPUはレイテンシを縮めず、大量スレッドという在庫で隠す。リトルの法則の $N$ を人力で用意している
SIMTは32スレッドに同じ命令を流す妥協。分岐が割れると時間は枝の和になり、最悪32倍
占有率はレジスタ・共有メモリ・ハード上限で決まる。ただし目的ではなく症状
Amdahlの $S(p) = 1/((1-f) + f/p)$ は着手前の足切り。上限は $1/(1-f)$

次は同じ並列化を複数台へ広げます。台数を増やしたとき新しく現れる制約は、演算でも帯域でもなく通信です。