WebGPU の新機能(Chrome 128)

François Beaufort
François Beaufort

サブグループのテスト

サブグループ機能を使用すると、SIMD レベルの並列処理が可能になり、グループ内のスレッドが通信して集約的な数学演算(16 個の数字の合計の計算など)を実行できます。これにより、非常に効率的なスレッド間データ共有が可能になります。

サブグループの提案の最小実装は、chrome://flags/#enable-unsafe-webgpu の「Unsafe WebGPU Support」フラグでローカルテストできます。

オリジン トライアルに登録して、実際のユーザーを対象にサイトのサブグループを試すこともできます。オリジン トライアルを使用するようにサイトを準備する手順については、オリジン トライアルのスタートガイドをご覧ください。オリジン トライアルは Chrome 128 ~ 131 で実施され、2025 年 2 月 19 日に終了します。テストの意図をご覧ください。

GPUAdapter"subgroups" 機能が使用可能な場合は、この機能を使用して GPUDevice をリクエストし、WGSL でサブグループをサポートし、minSubgroupSizemaxSubgroupSize の上限を確認します。

また、WGSL コードで enable subgroups; を使用して、この拡張機能を明示的に有効にする必要があります。有効にすると、次の追加機能が利用できるようになります。

  • subgroup_invocation_id: サブグループ内のスレッドのインデックスの組み込み値。
  • subgroup_size: サブグループ サイズ アクセスの組み込み値。
  • subgroupBallot(value): アクティブな呼び出しで value が true の場合は subgroup_invocation_id に対応するビットが 1、それ以外の場合は 0 であるビット フィールドのセットを返します。
  • subgroupBroadcast(value, id): id に一致する subgroup_invocation_id を持つ呼び出しから、サブグループ内のすべての呼び出しに value をブロードキャストします。注: id はコンパイル時の定数でなければなりません。

subgroupAddsubgroupAllsubgroupElectsubgroupShuffle などの組み込み関数は今後追加される予定です。問題 354738715 をご覧ください。

サブグループ オペレーションで f16 を許可するには、"subgroups""subgroups-f16""shader-f16" の各機能を使用して GPUDevice をリクエストし、WGSL コードで enable f16, subgroups, subgroups_f16; を使用して有効にします。

次のコード スニペットは、サブグループの可能性を探求するためのベースとなります。

const adapter = await navigator.gpu.requestAdapter();
if (!adapter.features.has("subgroups")) {
  throw new Error("Subgroups support is not available");
}
// Explicitly request subgroups support.
const device = await adapter.requestDevice({
  requiredFeatures: ["subgroups"],
});

const shaderModule = device.createShaderModule({ code: `
  enable subgroups;

  var<workgroup> wgmem : u32;

  @group(0) @binding(0)
  var<storage, read> inputs : array<u32>;

  @group(0) @binding(1)
  var<storage, read_write> output : array<u32>;

  @compute @workgroup_size(64)
  fn main(@builtin(subgroup_size) subgroupSize : u32,
          @builtin(subgroup_invocation_id) id : u32,
          @builtin(local_invocation_index) lid : u32) {
    // One thread per workgroup writes the value to workgroup memory.
    if (lid == 0) {
      wgmem = inputs[lid];
    }
    workgroupBarrier();
    var v = 0u;

    // One thread per subgroup reads the value from workgroup memory
    // and shares that value with every other thread in the subgroup
    // to reduce local memory bandwidth.
    if (id == 0) {
      v = wgmem;
    }
    v = subgroupBroadcast(v, 0);
    output[lid] = v;
  }`,
});

// Send the appropriate commands to the GPU...

ラインとポイントの深度バイアス設定のサポート終了

WebGPU の仕様変更により、レンダリング パイプラインのトポロジが線または点のタイプである場合に、depthBiasdepthBiasSlopeScaledepthBiasClamp をゼロ以外の値に設定すると検証エラーが発生します。デベロッパーがコードを更新するのに十分な時間を確保できるように、この検証が実施される前に DevTools コンソールに警告が表示されます。また、このような状況では値が強制的に 0 に設定されます。問題 352567424 をご覧ください。

未キャプチャ エラーの DevTools の警告を非表示にする(PreventDefault の場合)

DevTools コンソールで、uncapturederror のイベント リスナーが登録され、イベント リスナー コールバック内でイベント preventDefault() メソッドが呼び出されている場合、uncapturederror イベントに関する警告が表示されなくなりました。この動作は JavaScript のイベント処理と同じです。次の例と 問題 40263619 をご覧ください。

const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();

device.addEventListener("uncapturederror", (event) => {
  // Prevents browser warning to show up in the DevTools Console.
  event.preventDefault();

  // TODO: Handle event.error
});

WGSL はまずサンプリングを補間し、

WGSL interpolate 属性を使用すると、ユーザー定義の IO データ補間を管理できます。新たに補間サンプリング パラメータ first(デフォルト)と either が追加され、制御が強化されました。first ではプリミティブの最初の頂点の値が使用され、either では最初または最後の頂点のいずれかが許可されます。問題 340278447 をご覧ください。

Dawn の更新

非同期オペレーションを処理する Dawn の WGPUFuture の実装が完了しました。主なコンセプトには、オポチュニスティック イベント処理用の wgpuInstanceProcessEvents と、コールバック ロケーションを定義するための WGPUCallbackMode があります。WGPUFuture は、無限の存続期間を持つ 1 回限りのイベントを表します。wgpuInstanceWaitAny は、任意の Future の完了またはタイムアウトを待機します。問題 42240932 をご覧ください。

CompositeAlphaMode::Auto 値は Surface::GetCapabilities() によって報告されなくなりました。これは引き続き有効で、Surface::GetCapabilities().alphaMode[0] と同等です。問題 292 をご覧ください。

OpenGL バックエンドで、Present() 呼び出しごとに Y フリップ ブライトを使用する Surface がサポートされるようになりました。問題 344814083 をご覧ください。

Adapter::GetProperties() メソッドは非推奨となり、Adapter::GetInfo() の使用が推奨されています。

外部コントリビューターの Jaswant が、すべての CMake ファイルを書き直し、更新が容易になり、事前ビルドが可能になりました。CMake プロジェクトで Dawn を使用する方法については、クイックスタート ガイドをご覧ください。

以下に、主なハイライトをいくつかご紹介します。コミットの一覧(すべて網羅)をご覧ください。

WebGPU の新機能

WebGPU の新機能シリーズで取り上げられたすべての内容のリスト。

Chrome 130

Chrome 129

Chrome 128

Chrome 127

Chrome 126

Chrome 125

Chrome 124

Chrome 123

Chrome 122

Chrome 121

Chrome 120

Chrome 119

Chrome 118

Chrome 117

Chrome 116

Chrome 115

Chrome 114

Chrome 113