✍️ 필사 모드: ブラウザで本物のGPUコンピュート — WebGPUコンピュートシェーダーとWGSL実践ガイド 2026

プロローグ — ウェブがついにGPUを使いこなす

数年間、私たちは一つの嘘を受け入れて生きてきた。「ブラウザはGPUコンピュートができない」。WebGLにはシェーダーがあったが、それはグラフィックスパイプラインの中に閉じ込められたシェーダーだった。任意のデータを任意のワークグループに通し、結果を読み戻す — CUDAでは当たり前にできたこと — がウェブでは不可能だった。だからML推論はWebAssemblyでCPUに走らせるか、ONNX.jsのようにWebGLをコンピュートに見立てるトリックで搾り出した性能を使うしかなかった。

2023年にChrome 113がWebGPUをデスクトップ安定版に投入したところから風景が変わり始めた。2024年にFirefoxとSafariが続いた。そして2026年1月、Safari 26がmacOS TahoeとiOSで正式リリースされ、WebGPUは事実上Baselineに到達した。 Chrome・Edge・Firefox・Safariすべてがデフォルトで有効、グローバルカバレッジは約95%。

この記事はグラフィックスではない。コンピュートの話だ — 同じGPU、しばしば同じページに共存する、しかし別物のパイプライン。

• CPUで200msかかる行列積がGPUでは2msで終わる。
• WebLLMは量子化されたLlama 3 8Bをユーザーのノートパソコンで30+ tok/sで動かす。
• ガウシアンブラーを1000枚の画像にかける処理がwasm-SIMDで4秒なら、WebGPUなら0.4秒だ。

これらすべてを — サーバーコストゼロで — ブラウザ内で行う。これがコンピュートシェーダーが約束するものだ。

この記事はその約束を回収する方法を一気に整理する。

1章 · なぜブラウザでGPUコンピュートか

まず動機から。GPUコンピュートをブラウザで回すべき理由は3つだ。

1.1 サーバーコストゼロ

LLM推論をクラウドで動かすとH100が1時間あたり5–10ドル。ユーザーが100人ならコストはあっという間に積み上がる。同じ推論をユーザーデバイスのGPUに移せばサーバーコストはゼロだ。モデル重みは一度ダウンロードされてキャッシュされる。

1.2 プライバシー

医療画像、個人写真、私的なテキストを分類・要約・埋め込みする処理をサーバーに送らない。ブラウザのGPUで終わらせる。データはデバイスを離れない。

1.3 レイテンシ

サーバーまでの往復50–200msがなくなる。リアルタイム映像フィルタ、ライブML、インタラクティブシミュレーション — 往復がないこと自体が本質的に速い。

これらの理由は新しくない。新しいのはようやく道具が揃ったことだ。

2章 · WebGPU対WebGL — コンピュートの居場所

WebGL 2までのシェーダーはすべてグラフィックスパイプラインの中にいる。頂点シェーダーは頂点を、フラグメントシェーダーはピクセルを見る。任意のデータを任意に処理するには、データをテクスチャに偽装してフラグメントシェーダーに通す — いわゆる「GPGPUハック」を — するしかなかった。

WebGPUは違う。コンピュートシェーダーが一級市民だ。

核心は2行目だ。WebGPUはワークグループと共有メモリを与える。これがなければ私たちが「並列コンピュート」と呼ぶものの半分 — リダクション、スキャン、プレフィックス和 — は真似事しかできない。これらは協調アルゴリズムで、ワークグループ内の高速共有メモリを前提とする。

3章 · コンピュートモデル — ワークグループ、インボケーション、ディスパッチ

GPUは数千のALUを持つ。その数千に仕事を分配する方法がコンピュートモデルだ。

3.1 3つの階層

ディスパッチ (Dispatch)
  |
  +-- ワークグループ #0
  |     |
  |     +-- インボケーション (0,0,0)
  |     +-- インボケーション (1,0,0)
  |     +-- ...
  |     +-- インボケーション (63,0,0)
  +-- ワークグループ #1
  |     +-- ...
  +-- ワークグループ #N

• インボケーション (invocation): シェーダーコードを実行する最小単位。CUDAのthreadに相当。
• ワークグループ (workgroup): 同じSM(またはCU)で一緒に動くインボケーションの束。CUDAのblock。ワークグループ内では高速共有メモリとバリアで協調できる。
• ディスパッチ (dispatch): ワークグループの3Dグリッド。CUDAのgrid。

ワークグループサイズはシェーダーにコンパイル時定数として埋め込まれる。WGSLでは次のように書く(コードブロック内のみ — 散文では@compute @workgroup_size(64)のようにバッククォートで囲む)。

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  // ...
}

3.2 良いワークグループサイズは?

• 64または256が安全なデフォルト。 モバイルGPUもよく扱える。
• 大きすぎる(1024以上)と一部デバイスで拒否される。
• 小さすぎる(32未満)と占有率が落ちてGPUが飢える。

3.3 グローバルID

各インボケーションは自分のIDを知っている。global_invocation_idはディスパッチ全体内の3D座標。1D配列を処理するならgid.xだけ使えばよい。

4章 · WGSLクラッシュコース

WGSLはWebGPUのシェーダー言語だ。GLSLとHLSLの良いところを混ぜた — Rust風の — 文法を持つ。

4.1 型

// スカラー
var x: i32 = 42;
var y: u32 = 42u;
var z: f32 = 3.14;
var b: bool = true;

// ベクトル
var v: vec4<f32> = vec4<f32>(1.0, 2.0, 3.0, 4.0);
var i: vec3<i32> = vec3<i32>(1, 2, 3);

// 行列
var m: mat4x4<f32> = mat4x4<f32>(...);

// ランタイム配列
@group(0) @binding(0) var<storage, read_write> data: array<f32>;

4.2 ストレージバッファとバインドグループ

WGSLの入出力はバッファだ。グラフィックシェーダーと違って頂点アトリビュートもテクスチャもなく、GPUメモリに住む配列だけ。

// 読み取り専用入力
@group(0) @binding(0) var<storage, read> input: array<f32>;
// 読み書き出力
@group(0) @binding(1) var<storage, read_write> output: array<f32>;
// 小さな定数(1ディスパッチ中固定)
@group(0) @binding(2) var<uniform> params: Params;

groupとbindingはJS側のBindGroupLayoutと対応する。group 0 / binding 0はJSで設定した最初のバッファに繋がる、という具合。

4.3 ワークグループ共有メモリ

var<workgroup> shared_data: array<f32, 64>;

ワークグループ内の64個のインボケーションが同じ64要素配列を共有する。グローバル(ストレージ)メモリより約100倍速い。リダクション・スキャン・ブラーなどの協調アルゴリズムの中核資源。

4.4 バリア

ワークグループメモリに書いた値を他のインボケーションが見るには同期が必要だ。

workgroupBarrier();   // ワークグループ内の全インボケーションがここに到達するまで待つ
storageBarrier();     // ストレージバッファへの書き込みを全インボケーションに可視化

4.5 ビルトイン

よく使うビルトイン入力。

• global_invocation_id: ディスパッチ全体内の3D座標。
• local_invocation_id: ワークグループ内の3D座標。
• workgroup_id: ディスパッチ内のワークグループの3D座標。
• local_invocation_index: ワークグループ内の1Dインデックス(0..workgroup_size-1)。

5章 · 最初のコンピュートシェーダー — 並列和

理論は十分だ。実際に走らせる。長さNの配列を受けて全要素の和をGPUで計算する — クラシックな「最初のカーネル」 — 例。

5.1 WGSL — 2段リダクション

並列和は1ディスパッチで終わらせるのが難しい。ワークグループ間に共有メモリがないからだ。だから2段でやる。

1. 段階A: 各ワークグループが自分の範囲の部分和を計算してpartial_sums[wg_id]に書く。
2. 段階B: partial_sums配列をもう一度(またはホストで)減らして最終和を得る。

// reduce.wgsl — 段階A
const WG_SIZE: u32 = 64u;

@group(0) @binding(0) var<storage, read>       input:        array<f32>;
@group(0) @binding(1) var<storage, read_write> partial_sums: array<f32>;
@group(0) @binding(2) var<uniform>             params:       Params;

struct Params { n: u32 };

var<workgroup> sdata: array<f32, WG_SIZE>;

@compute @workgroup_size(WG_SIZE)
fn main(
  @builtin(global_invocation_id) gid: vec3<u32>,
  @builtin(local_invocation_id) lid: vec3<u32>,
  @builtin(workgroup_id) wid: vec3<u32>,
) {
  let i = gid.x;
  let tid = lid.x;

  // 1) 自分のグローバル要素を共有メモリにロード(範囲外なら0)
  var v: f32 = 0.0;
  if (i < params.n) { v = input[i]; }
  sdata[tid] = v;
  workgroupBarrier();

  // 2) ワークグループ内でツリー型に縮約
  var stride: u32 = WG_SIZE / 2u;
  loop {
    if (stride == 0u) { break; }
    if (tid < stride) {
      sdata[tid] = sdata[tid] + sdata[tid + stride];
    }
    workgroupBarrier();
    stride = stride / 2u;
  }

  // 3) 各ワークグループのインボケーション0だけが結果を書く
  if (tid == 0u) {
    partial_sums[wid.x] = sdata[0];
  }
}

3つの鍵となるパターン:

• グローバルからワークグループメモリ(sdata)へ一度だけロード。
• ツリーリダクション — strideを半分ずつ減らしてペアを足し合わせる。
• 各ステップ間にworkgroupBarrier() — メモリ整合性を保証。

5.2 JSドライバ

WebGPUのJS側は初見ではうるさい。だが部品は単純だ。

// reduce.ts
async function gpuSum(input: Float32Array): Promise<number> {
  // 1) アダプタとデバイス
  const adapter = await navigator.gpu.requestAdapter();
  if (!adapter) throw new Error('No WebGPU adapter');
  const device = await adapter.requestDevice();

  const N = input.length;
  const WG = 64;
  const numWorkgroups = Math.ceil(N / WG);

  // 2) バッファ群
  const inputBuf = device.createBuffer({
    size: input.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
  });
  device.queue.writeBuffer(inputBuf, 0, input);

  const partialBuf = device.createBuffer({
    size: numWorkgroups * 4,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
  });

  const uniformBuf = device.createBuffer({
    size: 16, // 16にパッド
    usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
  });
  device.queue.writeBuffer(uniformBuf, 0, new Uint32Array([N, 0, 0, 0]));

  // 3) シェーダとパイプライン
  const module = device.createShaderModule({ code: WGSL_SOURCE });
  const pipeline = device.createComputePipeline({
    layout: 'auto',
    compute: { module, entryPoint: 'main' },
  });

  // 4) バインドグループ
  const bindGroup = device.createBindGroup({
    layout: pipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: { buffer: inputBuf } },
      { binding: 1, resource: { buffer: partialBuf } },
      { binding: 2, resource: { buffer: uniformBuf } },
    ],
  });

  // 5) ディスパッチ
  const encoder = device.createCommandEncoder();
  const pass = encoder.beginComputePass();
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(numWorkgroups);
  pass.end();

  // 6) 結果をCPUに読むためのstagingバッファ
  const stagingBuf = device.createBuffer({
    size: numWorkgroups * 4,
    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
  });
  encoder.copyBufferToBuffer(partialBuf, 0, stagingBuf, 0, numWorkgroups * 4);
  device.queue.submit([encoder.finish()]);

  // 7) 読み取り
  await stagingBuf.mapAsync(GPUMapMode.READ);
  const partials = new Float32Array(stagingBuf.getMappedRange().slice(0));
  stagingBuf.unmap();

  // 8) 最終和はCPUで(numWorkgroupsが小さければ一度で十分)
  let total = 0;
  for (let i = 0; i < partials.length; i++) total += partials[i];
  return total;
}

初見では長いが3つに分かれる。

• セットアップ (1–4): アダプタ・デバイス・バッファ・パイプライン・バインドグループ。
• ディスパッチ (5): コマンドエンコード→キュー投入。
• 読み戻し (6–7): stagingバッファへコピー→mapAsync→読み取り。

5.3 性能 — 本当に速いのか

参考測定値 (M2 Pro, Chrome 127, N=8M float):

20倍速い。そしてNが大きくなるほど差が広がる。N=64Mでは約100倍。

6章 · 実際の応用 — コンピュートが本当に動いている場所

6.1 ブラウザLLM推論 — WebLLM

WebLLM(MLCチーム)が最もよく知られた例だ。Llama 3 8B、Qwen 2.5、Phi 3.5などのモデルを量子化してWebGPUコンピュートシェーダーで回す。M2 Proで4-bit Llama 3 8Bが30+ tok/s。サーバーコスト0円。

内部ではMLC-LLMのTVMがモデルグラフをWGSLコンピュートシェーダー列にコンパイルする。私たちが手で書くことは少ないが — アテンション・matmul・ソフトマックス — そのすべてが上で見たパターンのバリエーションだ。

6.2 ブラウザStable Diffusion

Web Stable DiffusionがSD 1.5とSDXL TurboをWebGPUで回す。7Bパラメータのモデルがユーザーのノートで1024x1024の画像を10–30秒で生成する。サーバーGPUとは比べられないが — サーバーコスト0、プライバシー100%、オフライン可。

6.3 GPUアクセラレートされたデータフレーム

GPU.jsの後継、WebDFのようなプロジェクトがpandas/Polars風のデータフレーム操作をWebGPUコンピュートで実装する。1000万行のgroup-by aggregationがCPUで400ms、GPUなら25ms。

6.4 並列画像フィルタ

最も即効性のある応用。ガウシアンブラー、エッジ検出、色空間変換を全ピクセルに同時適用する。ワークグループサイズを8x8または16x16に取って2Dステンシル処理。

@compute @workgroup_size(16, 16)
fn blur(@builtin(global_invocation_id) gid: vec3<u32>) {
  let px = gid.xy;
  // 3x3平均 — 本物のガウシアンは重みが違うだけ
  var sum: vec4<f32> = vec4<f32>(0.0);
  for (var dy: i32 = -1; dy <= 1; dy = dy + 1) {
    for (var dx: i32 = -1; dx <= 1; dx = dx + 1) {
      let p = vec2<i32>(i32(px.x) + dx, i32(px.y) + dy);
      sum = sum + textureLoad(input_tex, p, 0);
    }
  }
  textureStore(output_tex, vec2<i32>(px), sum / 9.0);
}

このようなフィルタはwasm-SIMDより10–50倍速い。

6.5 シミュレーション — N体、流体、布

物理シミュレーションはGPUの古典応用。1000–100000個の粒子の相互作用を毎フレーム解く必要がある。WebGPUコンピュートで60fpsで10000+粒子を処理可能。Babylon.jsやPlayCanvasが既に利用している。

7章 · 2026年のブラウザサポート現実

理論は良し、デモは派手だ。だが実際にデプロイするとぶつかる現実は別だ。

7.1 サポート状況

2026年5月時点(caniuse + WebKitリリースノート + Mozillaトラッカー総合):

グローバルカバレッジは約95%。残りの5%はwasm-SIMDまたは素のJSにフォールバック。

7.2 fp16 — まだ普遍的ではない

大型MLモデルのメモリフットプリントを半減するfp16(半精度)は — WGSLではf16拡張 — すべてのデバイスで有効ではない。 Chromeはデスクトップディスクリートで動くが、統合GPUやモバイルでは拒否されることもある。

const adapter = await navigator.gpu.requestAdapter();
const hasF16 = adapter.features.has('shader-f16');
const device = await adapter.requestDevice({
  requiredFeatures: hasF16 ? ['shader-f16'] : [],
});

フォールバックは2つ:

• f32に落とす — メモリ2倍、速度はほぼ同等。
• INT8量子化 — fp16より小さいが精度損失わずか。

7.3 ワークグループ限界

limits.maxComputeWorkgroupSizeXなどの限界はデバイスごとに違う。安全なデフォルト:

• workgroup_size: 64または256
• maxComputeInvocationsPerWorkgroup: 256(これ以上は絶対に埋め込まない)

7.4 メモリ限界

limits.maxStorageBufferBindingSizeがデフォルト128MB。大型モデルは複数バッファに分割する必要がある。WebLLMがモデル重みを複数バッファに分けるのもそのため。

7.5 コンパイル時間

WGSLシェーダーは初回使用時にコンパイルされる。大きなシェーダー(1000行のアテンションカーネル)は500ms–2s かかることがある。モデルロードと一緒にウォームアップ段階を設けるのが標準。

8章 · WebGPU + Wasm-SIMDハイブリッド

GPUがすべてに速いわけではない。小さな配列や分岐の多いコードはCPUのほうが速い。実際のパイプラインはしばしば両者を混ぜる。

8.1 いつGPU、いつCPU?

8.2 ハイブリッドパターン

WebLLMのようなMLパイプラインの典型的な流れ:

[トークナイザ]   <- wasm-SIMD(小さなコード、分岐多)
     |
     v
[埋め込み]       <- GPU(大きな行列積)
     |
     v
[アテンション]   <- GPU(matmul + softmax)
     |
     v
[FFN]            <- GPU(大きなmatmul)
     |
     v
[サンプラー]     <- CPUまたはGPU(小さなテンソル、分岐あり)
     |
     v
[デトークナイザ] <- wasm-SIMD

トークナイザ・デトークナイザのように小さく分岐の多いコードはCPU。行列積のように大きく均一なコードはGPU。GPU↔CPU間のデータ移動回数を最小化するのが核心 — 一度GPUに送ったら、できる限りそこで仕事を終わらせる。

8.3 WasmとWebGPUの通信

現状(2026年5月)、SharedArrayBufferを経由した直接共有はクロスオリジン分離を要求する。代替はstagingバッファを介したコピーで、PCIe帯域(数GB/s)に縛られる。可能な限りGPU上に留まることが答え。

9章 · 意思決定マトリックス

新機能を作るときどこで回すか、どう決めるか。

10章 · 限界と落とし穴

バラ色だけではない。2026年5月時点で知っておくべきこと。

10.1 初回ディスパッチは遅い

シェーダーコンパイル・パイプライン生成・バッファ割当はすべて初回使用時に起こる。短い1回限りの仕事ならGPUセットアップが本作業より高くつく。繰り返し使用または大きな仕事にのみ意味がある。

10.2 デバッグが難しい

GPUシェーダーにはconsole.logがない。WebGPUデバッグの現状は:

• Chrome DevToolsのWebGPUパネル — バッファ・バインドグループ・ディスパッチを検査。
• 出力バッファに値を書いてCPUで読んで確認(「printfデバッグ」)。
• WGSLのwgpu-inspectのようなツール — まだ未熟。

CUDAのcuda-gdbのような成熟したデバッガはない。

10.3 非決定論

浮動小数点和算順序はワークグループのスケジューリングで変わりうる。同じ入力にビット単位で同じ出力を期待しないこと。ML推論で同じモデル・同じプロンプトでも異なるトークンが出ることがある(sampling temperature 0でも)。

10.4 モバイルの揺らぎ

モバイルGPUはデスクトップより限界が厳しい。同じシェーダーがある携帯では動かないことがある。常に:

• adapter.limitsを読んでワークグループサイズとメモリ限界を確認。
• フォールバック経路を用意。
• 実機で検証(BrowserStack、実機両方)。

10.5 セキュリティ

WebGPUは大きな攻撃面だ。2023–2025年にかけていくつかの情報漏洩CVEが報告されている。ブラウザは強力に隔離するが — 機微なドメインではgpu-compute権限を慎重に検討。

10.6 シェーダーコードは保護されない

WGSLソースはクライアントに平文で渡る。アルゴリズムの企業秘密をそこに埋め込まないこと。

エピローグ — 2時間でコンピュートシェーダーを始める

開始チェックリスト

• WebGPUサポート確認: if (!navigator.gpu)ガード。
• アダプタ・デバイスは一度作って再利用(ページに1つ)。
• シェーダーモジュール・パイプラインをキャッシュ — 同じシェーダーを再コンパイルしない。
• バッファプーリング — 毎呼び出し新しいバッファを作らず再利用。
• stagingバッファとmapAsyncは非同期 — awaitを忘れない。
• ワークグループサイズは64または256から始める。
• 初回呼び出しはウォームアップ — 本番計測は5回目から。
• fp16はfeatureガード後に使用。

アンチパターン

• 毎呼び出しでデバイス/アダプタ再生成 — 高価で不要。
• GPU↔CPUピンポン — 反復ごとに結果をCPUへ読み戻さない。GPUに留まる。
• ワークグループサイズ1024+ — モバイルで拒否。
• 分岐だらけのシェーダー — GPU占有率を殺す。CPUに送る。
• fp16を当然視 — featureガードなしで爆発。
• 巨大単一シェーダー — コンパイルが長い。段階に分割。
• 結果を毎回mapAsync — async/awaitによるGPUパイプライン停止。バッチ化。

次回予告

• WGSL深掘り — アテンションカーネルを一から書く、flash attentionのワークグループトリック。
• WebGPUデバッグとプロファイリング実践 — Chrome DevTools、タイムスタンプクエリ、占有率測定。
• WebLLM内部 — TVMがモデルをWGSLにコンパイルするパイプラインの解剖。

参考 / References

• WebGPU標準 — W3C
• WGSL標準 — W3C
• MDN — WebGPU API
• WebGPU Fundamentals
• WebGPU Samples — webgpu.github.io
• WebLLM — MLC
• Web Stable Diffusion
• Chrome WebGPUリリースノート
• Safari 26リリースノート — WebKit
• Firefox WebGPUトラッカー — Bugzilla
• WebGPU Best Practices — Brandon Jones
• Surma — WebGPUコンピュート入門
• Babylon.js Compute Shaders文書