WWDC Quick Look 💓 By SwiftGGTeam
中文 EN 日本語
Unlock GPU computing with WebGPU

Unlock GPU computing with WebGPU

アプリサービスSafari と Web
元の動画を見る

ハイライト

WebGPU はブラウザ上で汎用 GPU 計算を実行できる現状唯一の API です。Apple プラットフォームでは Metal をバックエンドとして実装されており、Mac、iPhone、iPad、Vision Pro のすべてをカバーしています。

主要内容

これまで Web で 3D を扱う場合、開発者の選択肢は WebGL しかありませんでした。WebGL は compute shader をサポートしないため、ブラウザ内で物理シミュレーションやパーティクルシステム、ニューラルネットワーク推論を動かそうとすると、CPU でゴリ押しするか、ネイティブアプリへ処理を逃がすしかありません。さらに WebGL のステートマシン的な API には重い歴史的負債があり、draw call ごとに大量の境界チェックが走るため、パフォーマンスはネイティブに大きく劣ります。

WebGPU はこの問題をゼロから設計し直したものです。Safari チームの Mike が本セッションで示した結論は 3 点に整理できます。1 つ目は、WebGPU は WebGL でできることを 3D グラフィックスの領域でカバーしつつ、パフォーマンスと柔軟性のいずれにおいても上回るという点です (00:16)。2 つ目は、WebGPU はブラウザ上で汎用 GPU 計算を実行できる唯一の API であるという点 (00:28)。3 つ目は、WebGPU は Apple プラットフォームにおいて Metal をバックエンドに実装されており、Mac、iPhone、iPad、Vision Pro のすべてをサポートし、API の多くが Metal と 1 対 1 に対応しているため、Metal に慣れた開発者の学習コストはほぼゼロに近いという点です (00:38)。

API の設計はフラットな構造になっています。最上位に GPU と GPUAdapter があり、GPUDevice がほとんどの呼び出しのエントリポイントとなり、Metal の MTLDevice に対応します。リソース層には buffer、texture、sampler があり、エンコーダ (encoder) がコマンドを発行、pipeline がリソースの解釈方法を記述、bind group が関連リソースをまとめ、内部的には Metal の argument buffer にマッピングされます。シェーダー言語 WGSL は Web のセキュリティモデルに合わせてゼロから設計され、vertex、fragment、compute の 3 種類のプログラムをサポートします。セッションでは 10 万三角形 + パーティクル物理シミュレーションのサンプルを使って、パイプライン全体を一気通貫で説明しています。

詳細

デバイス初期化と f16 拡張

WebGPU アプリケーションはすべて GPUDevice の取得から始まります。Mike は (05:14) で shader-f16 拡張を含むデバイス生成のコードを示しています。

const canvas = document.querySelector('canvas');
const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice({
  requiredFeatures: ['shader-f16'],
});

const context = canvas.getContext('webgpu');
context.configure({ device, format: 'bgra8unorm' });

ポイントは次のとおりです。

  • navigator.gpu.requestAdapter() は現在のシステムで利用可能な GPU アダプタを返します。物理デバイスを選択する操作に相当します。
  • requestDevice がほとんどの後続 API のエントリポイントで、Metal の MTLDevice に相当します。
  • requiredFeatures: ['shader-f16'] は半精度浮動小数点を有効化するもので、VRAM 帯域を半分に削減できますが、すべてのプラットフォームで対応しているわけではないため、利用前に可用性を確認する必要があります (05:39)。
  • context.configure は canvas を GPU が書き込み可能なメモリと結び付けるもので、これを行ってはじめてレンダリング結果を画面に表示できます。

buffer と texture の生成

リソース層の中核となるオブジェクトは buffer と texture の 2 つで、それぞれ MTLBuffer と MTLTexture に対応します (06:24)。

const particleBuffer = device.createBuffer({
  size: byteLength,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
});
device.queue.writeBuffer(particleBuffer, 0, particleData);

const texture = device.createTexture({
  size: [width, height],
  format: 'rgba8unorm',
  usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST,
});
device.queue.copyExternalImageToTexture(
  { source: imageBitmap },
  { texture },
  [width, height],
);

ポイントは次のとおりです。

  • usage はビットマスクで、このメモリの用途を WebGPU に伝えます。明示的に宣言しておくことで、API の複雑さを増やさずにデータ競合を回避できます (07:00)。
  • device.queue.writeBuffer は buffer、オフセット、JS の ArrayBuffer を引数に取り、JS 側のデータを GPU から見えるメモリへコピーします。
  • texture は copyExternalImageToTexture を使って ImageBitmap からアップロードし、テクスチャ画像として用いるのが一般的です。
  • texture は 1D、2D、2D 配列、cube map (2D テクスチャ 6 枚)、3D の 5 種類の形態をサポートします (07:44)。

WGSL で compute shader を書く

WebGL に compute shader が存在しなかったのに対し、WebGPU では使えるようになった点が最大の能力アップです (12:17)。パーティクル物理を扱う compute shader の骨格は次のようになります (12:26)。

enable f16;

@group(0) @binding(0) var<storage, read_write> particles: array<Particle>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  var p = particles[gid.x];
  p.velocity += gravity * deltaTime;
  p.position += p.velocity * deltaTime;
  particles[gid.x] = p;
}

ポイントは次のとおりです。

  • enable f16 でシェーダー側でも半精度型を有効化します。これは device 側の shader-f16 拡張と必ずセットで使う必要があります (15:24)。
  • @workgroup_size(64) は各 workgroup のグリッドサイズを定義し、compute shader の並列度の粒度を決定します。
  • @builtin(global_invocation_id) は WGSL の組み込み変数で、現在のスレッドが dispatch グリッド全体で占める位置を返してくれるため、JS からパラメータとして渡す必要はありません (12:34)。
  • var<storage, read_write> は読み書き可能なストレージ buffer を宣言します。読み取り専用で十分な場合は read を使うべきで、こうすることで WebGPU が一部の検証をスキップできるようになります (16:38)。

Render bundle で描画コマンドをキャッシュする

WebGPU はフレームごとに読み書き境界をすべて検証する仕組みになっており、静的なシーンではこのオーバーヘッドが純粋な無駄になります。Render bundle はこの検証を 1 度の生成時へ前倒しできます (17:04)。

const encoder = device.createRenderBundleEncoder({
  colorFormats: ['bgra8unorm'],
});
encoder.setPipeline(pipeline);
encoder.setBindGroup(0, bindGroup);
encoder.setVertexBuffer(0, vertexBuffer);
encoder.draw(vertexCount);
const bundle = encoder.finish();

// フレームごとに再生
renderPass.executeBundles([bundle]);

ポイントは次のとおりです。

  • createRenderBundleEncoder はあくまでレコーダで、finish() を呼び出してはじめて再生可能な bundle が生成されます。
  • executeBundles([bundle]) でコマンド全体を一度に再生し、フレームごとの検証をスキップします。
  • 内部的には Metal の indirect command buffer にマッピングされており、性能はネイティブ Metal と同等のメリットが得られます (18:03)。

Dynamic offset で bind group をマージする

各 bind group は 1 つの MTLBuffer に対応するため、数が増えると Metal オブジェクトが膨らんでしまいます。dynamic offset を使えば、1 つの bind group で同じレイアウトを使い回しつつ、実行時にオフセットだけを切り替えることが可能です (19:49)。

const layout = device.createBindGroupLayout({
  entries: [{
    binding: 0,
    visibility: GPUShaderStage.VERTEX,
    buffer: { type: 'uniform', hasDynamicOffset: true },
  }],
});

const bindGroup = device.createBindGroup({
  layout,
  entries: [{ binding: 0, resource: { buffer: bigBuffer, size: 64 } }],
});

renderPass.setBindGroup(0, bindGroup, [offset]);

ポイントは次のとおりです。

  • hasDynamicOffset: true を使うには、shader の auto layout ではなく、自前で定義したカスタム layout で生成する必要があります (20:07)。
  • setBindGroup の第 3 引数に渡すのが動的オフセットの配列で、動的 buffer 1 つにつきオフセットを 1 つ指定します。
  • Mike が示した比較例では、64 バイトの bind group を 10 個用意するよりも、640 バイトの buffer 1 つ + オフセット 10 個に集約することで Metal buffer を 9 個削減できます (20:43)。

Command buffer と pass の数

Command buffer の境界では、オンチップの高速メモリと統合メモリ間で同期を取る必要があり、そのオーバーヘッドは無視できません。基本的には更新ループごとに 1 つの command buffer にまとめておき、データを統合メモリへ書き戻す必要があるときだけ分割するのが推奨です (18:28)。Pass 同士には同期は不要ですが、各 pass が帯域を消費するため、できる限りまとめるべきです。Apple GPU は tile-based deferred renderer であるため、pass を集約する利点はデスクトップ GPU よりも顕著に現れます (19:15)。

重要ポイント

1. 新しい Web プロジェクトでは WebGPU に直接賭け、WebGL は飛ばす

なぜ取り組む価値があるのか: WebGPU は Apple の全プラットフォームで Metal を介して実装されており、性能上限が WebGL とは桁違いです。さらに compute shader を実行できるのは WebGPU だけです。WebGL の能力上限では、LLM 推理、物理シミュレーション、複雑なパーティクルといったアプリケーションを支えるのは無理があります。

どう始めるか: three.js や Babylon.js のようにすでに WebGPU バックエンドをサポートしているライブラリから始めて、renderer を WebGPURenderer に切り替えます。低レベルから書きたい場合は navigator.gpu.requestAdapter() を直接使い、WebGPU Samples を参照すると良いでしょう。

2. ブラウザ上の推論と物理計算を compute shader へ移植する

なぜ取り組む価値があるのか: ブラウザ内で transformer モデル、剛体シミュレーション、流体シミュレーションを動かそうとすると、これまでは WASM SIMD で頑張るか、諦めるかの二択でした。compute shader を使えばこれらの処理が GPU の並列度をフルに活かせるようになり、ブラウザを離れずに完結するためデプロイコストはほぼゼロです。

どう始めるか: まずは Transformers.js (WebGPU バックエンド対応済み) のオンデバイス推論を動かしてみるのがおすすめです。自前のモデルを書く場合は WGSL の @compute @workgroup_size と global_invocation_id を参考に、最小構成の kernel から作っていくと良いでしょう。

3. 静的・準静的なシーンはすべて render bundle で扱う

なぜ取り組む価値があるのか: UI 要素や更新頻度の低い 3D シーンに対してフレームごとに検証を走らせるのは純粋な無駄です。Render bundle なら検証を 1 度で済ませられ、性能はネイティブ Metal に肉薄します。

どう始めるか: 既存の render pass からフレームごとに繰り返されているコマンドを抜き出し、createRenderBundleEncoder で一度だけ録画しておき、メインループは executeBundles に置き換えましょう。動的な部分は引き続き通常の pass で処理します。

4. iOS / visionOS では f16 + dynamic offset でメモリを抑えるのを基本にする

なぜ取り組む価値があるのか: Mike も 15:01 で触れているとおり、iOS と visionOS はメモリに非常に敏感で、上限を超えるとシステムにプロセスを kill されてしまいます。f16 にするだけで帯域は半分になり、dynamic offset で Metal オブジェクト数を減らせるため、両者を組み合わせるのがプロセスの寿命を延ばす最も直接的な手段です。

どう始めるか: requestDevice で requiredFeatures: ['shader-f16'] を指定し、WGSL 側でも enable f16 を有効化したうえで、まずは重みや頂点といった大きな読み取り専用データから f16 に切り替えていきます。bind group はすべて dynamic offset を使い、同種のオブジェクトをまとめて大きな buffer に集約しましょう。

関連セッション

コメント

GitHub Issues · utterances