Kotonia Articles

LLMが2%に落ちた。GPU同居の天井を、3回測り直した

LLMのデコードは帯域律速、拡散モデルは計算律速。使う資源が違うのだから同じGPUに同居できるはず——という仮説を、CUDA MPSとSMパーティショニングで3回測り直した記録。結論は「空間ではなく時間」だった。

著者清水真二2026-07-318分で読める

#gpu#cuda#vllm#blackwell#llm

他の言語英語中国語

はじめに

手元のマシンにGPUが2枚あって、そこに音声対話・リップシンクアバター・LLM・画像生成・動画生成を全部載せている。もう1枚届くので、どのワークロードをどのカードに置くかを設計していた。

その過程で、筋の良さそうな仮説が出た。

LLMのデコードはメモリ帯域律速で、拡散モデルは計算律速。使う資源が違うのだから、同じカードに載せても重ならないのではないか？

NVIDIAにはMPS（Multi-Process Service）という、まさにこれを可能にしそうな機能がある。試した。

結論から言うと、仮説は間違っていた。そして「間違っている」と結論した過程も、2回間違っていた。 3回測り直して、ようやく説明のつく形になった。その3回ぶんを順に書く。

第1幕: 空間の仮説

2枚のカードのルーフラインを測る

議論の土台がないので測った。RTX PRO 6000 Blackwell Max-Q（96GB）と RTX PRO 4000 Blackwell（24GB）。

	6000 Max-Q	4000	比
SM数	188	70	2.69x
メモリ帯域	1469 GB/s	553 GB/s	2.66x
bf16 GEMM (sustained)	228 TFLOPS	103 TFLOPS	2.22x
リッジ点	155.5 FLOP/byte	185.9 FLOP/byte	—
4秒でのスロットル	-2.0%	-0.9%	—

両カードともメモリクロック14001MHz / SMクロック3090MHzで同一だった。差はバス幅・SM数・電力枠だけ。帯域比（2.66x）が計算比（2.22x）より大きく、6000のほうがリッジ点が低い＝相対的に帯域が潤沢なカードということになる。

Max-Qの300W制限が計算を削っているのでは、と疑っていたが、4秒のGEMMでスロットルは-2.0%しか出なかった。

仮説の根拠

リッジ点は「この演算強度を境に、帯域律速と計算律速が入れ替わる」点だ。6000なら155.5 FLOP/byte。

LLMのデコードはこの遥か下にいる。 トークンを1個出すには全レイヤーの全重み行列を頭から読み、それぞれに「今のトークン1個ぶんの細いベクトル」を掛ける。重みは1回読んで1回使って捨てる。再利用がゼロなので、重みの格納精度をb byte/paramとして演算強度は2B/b（Bはバッチ）にしかならない。NVFP4（0.5 byte/param）でバッチ1なら4 FLOP/byte。リッジ点の1/39。

測った帯域を27BのNVFP4モデル（約13.5GB）に当てると、

13.5 GB / 1469 GB/s = 9.2 ms/token  →  理論上限 およそ109 tok/s

LLMの速度は計算力ではなく「VRAMから重みをどれだけ速く吸い出せるか」で決まっている。テンソルコアはほぼ遊んでいる。

拡散モデルは逆側にいる。 2048²の画像生成は読んだデータを何千回も使い回す。演算強度は数百〜数千。帯域はガラガラでテンソルコアだけが張り付く。

資源の使う場所が綺麗に分かれている。だから重なるはず——という仮説だった。

その前に: GPUはそもそも別プロセスのカーネルを同時に走らせない

GPUは「1プロセス = 1 CUDAコンテキスト」で動き、別コンテキストのカーネルは同じSMの上で同時には走らない。GPUがコンテキストごと切り替えて順番に処理する。時分割だ。

MPSなし — 時分割。常にどちらか片方だけが走る
時間 →
LLM       ████░░░░████░░░░████░░░░
画像生成   ░░░░████░░░░████░░░░████
SM使用率   低  高  低  高  低  高    ← LLMのターンはテンソルコアが遊ぶ
帯域使用率 高  低  高  低  高  低    ← 画像のターンは帯域が遊ぶ

MPSはこれを変える。デーモンが全プロセスのGPU作業を1つの共有コンテキストに束ねる代理人になるので、別プロセス同士のカーネルが本当に同時に走れるようになる（はずだった）。

計器を作る

本物同士をいきなり戦わせると何が効いたのか分からない。本物のカーネルではなく、本物の「ルーフライン上の位置」だけを再現する代理を2つ作った。

帯域側（LLM decodeの代理）——巨大な行列 × 細いベクトル。1回で537MBを読んで537 MFLOPしか計算しない。演算強度1.0 FLOP/byte。

w = torch.randn(16384, 16384, dtype=torch.bfloat16, device="cuda")  # 537 MB
x = torch.randn(16384, 1, dtype=torch.bfloat16, device="cuda")      # バッチ1
while ...:
    torch.mm(w, x, out=y)

計算側（拡散モデルの代理）——正方行列同士。1回で1.6GB触って8.8 TFLOP計算する。演算強度5461 FLOP/byte、リッジ点の35倍上。

a = torch.randn(16384, 16384, dtype=torch.bfloat16, device="cuda")
while ...:
    torch.mm(a, b, out=c)

素朴に2つ同時起動すると「片方が先に終わって残りが独占した時間」が混ざるので、共通のUNIX時刻を両プロセスに渡し、ウォームアップ後にバリアで待ち合わせて完全に同じ5秒窓で計測する。

そして計器の本体は「正規化した和」だ。 各条件のスループットを単独実行時の値で割り、2つを足す。

和 ≈ 1.0  →  ゼロサム。片方の得は他方の損 = 重なっていない
和 >  1.0  →  本当に重なっている。遊んでいた資源が埋まった

「重なっているのか、順番待ちなのか」を1個の数字に潰す。

結果: 時分割はゼロサム、MPSは帯域側を餓死させる

単独基準は帯域側1435 GB/s、計算側232 TFLOPS。

	帯域側	計算側	和
MPSなし	44%	51%	96%
MPS / 上限なし	2%	98%	100%

MPSなしはぴったりゼロサム。理論通りで、計器が壊れていないことは確認できた。

そしてMPSを入れると、帯域側が時分割より21倍悪化した。

なぜか: 「計算を使わない」と「SMが要らない」は別

ここが仮説の穴だった。メモリを読むにも、スレッドがSMに居座ってロード命令を発行する必要がある。 「テンソルコアを使わない」ことと「SMが要らない」ことは全く別だ。

計算側のGEMMは1発37.9msの化け物カーネルで、188個のSMを埋め尽くすthread blockを積む。しかもblock粒度でのプリエンプションが無いので、一度載ったblockは最後まで走る。帯域側のGEMVはその列の後ろに並ぶだけで、そもそも発射できない。

時分割のときは少なくとも自分の順番が回ってきた。MPSにすると「並んで走れる」代わりに「席が空かない」状態になる。

SM上限を掛けると戻る。ただし常時課税される

MPSにはCUDA_MPS_ACTIVE_THREAD_PERCENTAGEという、クライアントごとにSMの取り分を縛るノブがある。計算側を絞ると帯域側が戻ってくる。

計算側のcap	帯域側	計算側	和
上限なし	2%	98%	100%
50%	61%	43%	104%
25%	85%	22%	107%
10%	93%	7%	100%

相手を走らせずに単独で測ると、capは競合時だけの調停ではなく固定税だと分かった。cap 25%なら誰も居なくても46%しか出ない。

cap	帯域側（単独）	計算側（単独）
100%	100%	100%
50%	93%	75%
25%	61%	46%

ここで重要な非対称が見える。帯域律速側はSMを半分に削っても93%出る。 半分あれば帯域はほぼ飽和する。

両方にcapを掛けても、SMは分割されない

単独capの実測から「もし本当にSMが分割されるなら」の予測が立つ。帯域側50%（単独93%）＋計算側50%（単独75%）なら和は168%のはずだ。

設定	分割されるなら	実測
帯域50 / 計算50	168%	107%
帯域50 / 計算25	139%	102%
帯域75 / 計算25	~143%	109%
帯域25 / 計算75	~149%	99%

まったく届かない。どう振っても和は99〜109%に貼りついて、天井が約107%から動かない。capは「使ってよいSM数の上限」であって「取り置く」ではない。取り上げることはできても、保証することはできない。

ここで結論を出しかけた

1枚のGPUを2つの飽和ワークロードで分け合う方向には、ハードウェア的な天井がある。時分割はゼロサム、MPSはどう設定しても約107%、MIGはこのカード世代では非対応。ソフトでは越えられない。

きれいな結論だった。数字も揃っていた。そして間違っていた。

第2幕: 計器を疑う

引っかかったのは「capは上限であって予約ではない」という部分だった。なら、本当にdisjointなSM集合を作ったらどうなるのか。

CUDAには%smidという、自分がどのSMで走っているかを読める特殊レジスタがある。担当外のSMに載ったblockを即returnさせれば、ライブラリに頼らず本物の分割が作れる。

__device__ __forceinline__ unsigned smid() {
  unsigned r; asm volatile("mov.u32 %0, %%smid;" : "=r"(r)); return r;
}

__global__ void bw_pass(...) {
  unsigned s = smid();
  if (s < lo || s >= hi) return;    // 担当外のSMは即降りる
  ...
}

第1幕の失敗も踏まえて作り直した。永続ブロック（締切まで居座る）をやめて実カーネル同様にlaunchを繰り返す形にし、バッファを16GiBにしてL2（128MiB）を確実に溢れさせ、masked blockが抜けてもデータを取りこぼさないようアトミックなワークキューにした。

結果: 分割は効かない。しかし和が139%あった

帯域側 単独 (全SM)          1650 GB/s
帯域側 単独 (下半分のSM)    1651 GB/s   ← 半分のSMで帯域を完全に飽和 (100%)
計算側 単独 (上半分のSM)    53.9 TFLOPS (56%)
  → 分割が完全に効くなら 和 157% のはず

同時 / SM共有               和 139%
同時 / SM分割               和 138%   ← 変わらない

明示的にSMを分割しても和は上がらなかった。 帯域律速側が半分のSMで既に飽和していたので、分割の余地が最初から無かった。ハードの貪欲スケジューラは、このペアについては既にほぼ最適だった。

だが目を引いたのはそこではない。和が139%ある。 MPSの天井107%とも、時分割の96%とも全く違う。

同じカーネルでプロセスを分けてみる

第1幕はcuBLAS、第2幕は自作カーネル。カーネルが違えば比較できない。同じカーネルのまま構成だけ変えて測り直した。

構成	帯域側	計算側	和
2プロセス / MPS OFF	48%	48%	96%
2プロセス / MPS ON	66%	66%	132%
1プロセス2ストリーム / SM共有	91%	48%	139%
1プロセス2ストリーム / SM分割	95%	44%	138%

MPSは96% → 132%で普通に効いていた。 第1幕の「天井は+7%」は、一般命題としては嘘だった。

犯人はカーネル粒度だった

差は2つあった。

1発の長さが7倍違う。

cuBLAS GEMM: 37.9 ms/発
自作カーネル: 5.2 ms/発

ブロック粒度のプリエンプションが無い以上、「空くまで待つ」しか手が無い。この7倍がそのまま割り込みやすさになる。

占有率に隙間を残していた。 実測すると:

SM あたり最大スレッド数    : 1536
自作カーネル              : レジスタ 20/thread、shared mem 0 B
  → 理論上 6 block/SM (占有率100%) まで詰められる
  → でも実際は 4 block/SM しか launch していなかった
  → SM のスレッドスロットが 33% 空いたまま

その空いた33%が、まさに相手のブロックが入り込んだ場所だった。cuBLASは逆で、大きなshared memoryタイルと大量のレジスタで占有率を使い切るようにチューニングされているので、他人の入る隙が無い。

つまり僕の代理は、拡散モデルの代理としては軽すぎた。 演算強度というルーフライン上の位置は合っていたのに、SMの占有の仕方が全く違った。ルーフラインだけ合わせても代理にならなかった。

第3幕: 実物で答え合わせ

代理が信用できないなら、実物で測るしかない。ただし今回は先に予測を書いてから測った。

HiDreamの重い層（attentionとlinear）は結局cuBLAS / cuBLASLtを呼んでいて、それはまさに「占有率を使い切る、長い」カーネルだ。だから悲観側に寄るはず。ただし拡散の1ステップは巨大GEMM1発ではなく、layernorm・要素演算・softmax・小さい射影が数百個並んだ列で、その間は占有率が落ちる。だから107%（cuBLAS相当）と132%（細粒度相当）の間、下端寄りに落ちる。

本番からThinkCap-27B（NVFP4, vLLM）とHiDream-O1-Imageを隔離し、GPUを空にして測った。画像は本番と同じ2048²のまま、ステップ数だけ50→10に落としている。ステップは同じカーネルの反復回数なので、粒度を保ったまま実験時間だけ縮められる。

結果

基準は LLM 71.83 tok/s（TTFT 62.0ms）、画像 8.12秒/枚。

条件	LLM		TTFT	画像		和
MPS OFF（現状）	33.18 tok/s	46.2%	84.5 ms	15.05 s/枚	54.0%	100.1%
MPS ON	18.53	25.8%	126.1 ms	9.13	88.9%	114.7%
MPS ON + 画像 SM 25%	52.80	73.5%	67.2 ms	36.86	22.0%	95.5%

114.7%。予測した範囲の下端に着地した。

ちなみに基準の71.83 tok/sは、第1幕で帯域から予測した上限109 tok/sの66%にあたる。MTPのオーバーヘッドとKVキャッシュ読み出しを考えれば妥当な線で、ルーフラインのモデル自体は生きている。

三者三様で、しかも「最良」が指標によって全部違う

総和が最大なのはMPS ON（114.7%）。でもLLMは25.8%まで落ちる
総和が最小なのはMPS + cap（95.5%、時分割より下）。でもLLMは73.5%で最良
時分割は全部の中間

会話体験を決めるのはLLMのトークン速度とTTFTだ。そこだけ見ると:

今 (時分割)        : 46.2% / TTFT 84.5ms
MPS + 画像 cap 25% : 73.5% / TTFT 67.2ms   ← 単独時の 62.0ms にほぼ復帰

画像生成中の会話速度が1.6倍、TTFTはほぼ無負荷時まで戻る。代償は画像が15.1秒 → 36.9秒（2.4倍遅い）。画像は非同期でキューに積めるので、トレードとしては明確に得だ。

総和が最小の設定が、プロダクトとしては最良になる。 スループットの総和は、この問題における通貨として間違っていた。

結論: 空間ではなく時間

3回測って、最後に残ったのはこれだ。

空間の議論——どの資源を食うか、どのSMに置くか——は、最後まで一度も効かなかった。

出発点だったルーフライン分析は筋が通っていたし、実際に正しい。ただし単独で走らせたときの速度については。同居の可否については、ほぼ何も予測していなかった。データは何度もそう言っていた。

帯域律速側は半分のSMで帯域を100%飽和させた（＝空間を分ける余地が最初から無い）
明示的にSMをdisjointに分けても和は139% → 138%（＝空間的分離に価値が無い）
一方で、カーネルの長さを37.9msから5.2msに変えただけで107% → 132%に動いた

効いていたのは一貫して「どれだけ割り込めるか」だった。そしてそれは2つの掛け算だ。

	問い	決めるもの
回転	いつ席が空くか	カーネル1発の長さ（block粒度のプリエンプションが無いので、終わるまで待つしかない）
空席	空いたとき入れる余地があるか	隣人が使い残した占有率

どちらも入場できるかの話で、入った後にどの演算ユニットを使うかとは無関係だ。

面白いのは、唯一効いた空間的なノブ（SM cap）も、実は時間軸で効いていたことだ。あれは誰かに専用のSMを与えているのではなく、貪欲なクライアントの足跡を小さくすることで席が空く頻度を上げていた。だから純粋な空間分割は無価値なのに、上限としてのcapは効く。

実務的な帰結

他人のコードでは打つ手が無い。 カーネルを短くするのも占有率に余裕を残すのも、自分で書いたカーネルなら設計できる。だがvLLMとcuBLASの中身は触れない。サードパーティのスタックを同居させる問題は、原理的に物理分離にしか解が無い。

そして自分のプロダクトの話になる。うちはリアルタイムの音声ロールプレイで、会話ループは「発話検出 → 音声認識 → LLM → 音声合成 → リップシンクアバター」だ。アバターのフレーム生成にはハードな締切がある。

このモデルで見ると、アバターが同居で壊れる理由がはっきりする。アバターは400msの締切ごとに、encode → 9回のdenoise step → decodeという依存鎖を持っている。鎖の各リンクが入場の抽選を個別に受ける。 1回のスループット低下ではなく、抽選を11回連続で通す必要がある。だから平均で1.3倍遅い程度でも、最悪ブロックは締切を割る。

答えは物理分離になる。リアルタイム経路（音声とアバター）を専用カードに隔離し、LLMと画像・動画生成は別のカードへ。届く3枚目は、この線引きを完成させるために使う。

MPS + capは入れない。LLM 73.5%は魅力的だが、デーモンが単一障害点になり、vLLMをMPS下で長時間走らせた実績も無く、そして何よりカードが来れば問題自体が消える。「1枚に同居させざるを得なくなったときの手札」として記録しておく。

3回測り直して学んだこと

計器の妥当性検算は「壊れていない」ことしか保証しない。 第1幕でMPSなしがぴったり96%（ゼロサム）になったのを見て、僕は計器を信用した。実際その計器は壊れていなかった。ただ代表していなかった。この2つは別物だ
代理を作るとき、ルーフライン上の位置だけ合わせても足りない。 占有の仕方まで合わせないと、同居の実験には使えない
予測を先に書いてから測ると、当たったときに「モデルに予測力がある」と言える。 第3幕で114.7%を当てられたので、次に別のワークロードを載せるときは、実験する前に机上で見当がつく。これが3回ぶんの本当の収穫だった

この実験の限界

測ったのは性格が両極な2者のペアだけ。3者以上や、CPU側の同期が重いワークロードで同じ形になるかは未検証
動画生成（LTX-2）は「88秒のうち半分近くはディスクからのロードでGPUがアイドル」なので、このモデルだとむしろ良い隣人という予測になる。測っていない
cap付き画像の「単独時」の速度を測っていない。 合成負荷ではcapは無条件税だったので、実物でも同じなら深夜の無負荷時も画像が遅いままになる。デプロイを検討するなら必須の追試
5秒〜45秒の窓なので、熱平衡は見ていない