WebGPU チュートリアル: Web 上のコンピューティング、頂点、およびフラグメント シェーダー

WebGPU は、最先端の GPU コンピューティング機能を Web にもたらし、共有コード ベースを使用してすべての消費者プラットフォームに利益をもたらすことを約束するグローバル テクノロジーです。

その前身である WebGL は強力ですが、コンピューティング シェーダー機能が著しく不足しており、アプリケーションの範囲が制限されています。

WGSL (WebGPU Shader/Compute Language) は、Rust や GLSL などの分野のベスト プラクティスを活用しています。

WebGPU の使い方を学んでいると、ドキュメントにいくつかのギャップがあることに気づきました。計算シェーダーを使用して頂点シェーダーとフラグメント シェーダーのデータを計算するための簡単な出発点を見つけたいと思っていました。

このチュートリアルのすべてのコードの単一ファイル HTML は、https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb にあります。詳細な内訳については、以下をお読みください。

これは、私のドメインで実行されるこの HTML のシングルクリックのデモです: https://www.php.cn/link/bed827b4857bf056d05980661990ccdc Chrome や Edge などの WebGPU ベースのブラウザhttps://www.php.cn/link/bae00fb8b4115786ba5dbbb67b9b177a)。

詳細設定

これは粒子シミュレーションです。時間の経過とともにタイムステップで発生します。

時間は JS/CPU で追跡され、(float) 均一として GPU に渡されます。

パーティクル データは完全に GPU 上で管理されますが、CPU と対話し続けるため、メモリの割り当てと初期値の設定が可能になります。データを CPU に読み戻すことも可能ですが、このチュートリアルでは省略します。

このセットアップの魅力は、各パーティクルが他のすべてのパーティクルと並行して更新され、ブラウザでの驚異的な計算速度とレンダリング速度が可能になることです (並列化により GPU のコア数が最大化されます。パーティクルの数は次のように割ることができます)。コアごとの更新ステップごとの真のサイクル数を取得するには、コアの数を使用します）。

バインド

WebGPU が CPU と GPU 間のデータ交換に使用するメカニズムはバインドです。WebGPU バッファーを使用して、JS 配列 (Float32Array など) を WGSL のメモリ位置に「バインド」できます。 WGSL メモリの場所は、グループ番号とバインディング番号という 2 つの整数によって識別されます。

私たちのケースでは、コンピューティング シェーダーと頂点シェーダーの両方が、時間とパーティクルの位置という 2 つのデータ バインディングに依存しています。

時間 - 制服

コンピューティング シェーダー (https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L43) と頂点シェーダーには均一な定義が存在します。 (https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L69) 中 - シェーダー更新位置を計算し、頂点シェーダーは時間に基づいて色を更新します。

コンピューティング シェーダーから始めて、JS と WGSL のバインディング セットアップを見てみましょう。

<code>const computeBindGroup = device.createBindGroup({
  /*
    参见 computePipeline 定义，网址为
    https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L102

    它允许将 JS 字符串与 WGSL 代码链接到 WebGPU
  */
  layout: computePipeline.getBindGroupLayout(0), // 组号 0
  entries: [{
    // 时间绑定在绑定号 0
    binding: 0,
    resource: {
      /*
      作为参考，缓冲区声明为：

      const timeBuffer = device.createBuffer({
        size: Float32Array.BYTES_PER_ELEMENT,
        usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST})
      })

      https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L129
      */
      buffer: timeBuffer
    }
  },
  {
    // 粒子位置数据在绑定号 1（仍在组 0）
    binding: 1,
    resource: {
      buffer: particleBuffer
    }
  }]
});</code>

およびコンピューティング シェーダー内の対応する宣言

<code>// 来自计算着色器 - 顶点着色器中也有类似的声明
@group(0) @binding(0) var<uniform> t: f32;
@group(0) @binding(1) var<storage read_write=""> particles : array<particle>;
</particle></storage></uniform></code>

重要なのは、JS と WGSL のグループ番号とバインディング番号を一致させることによって、JS 側の timeBuffer を WGSL にバインドしていることです。

これにより、JS から変数の値を制御できるようになります:

<code>/* 数组中只需要 1 个元素，因为时间是单个浮点值 */
const timeJs = new Float32Array(1)
let t = 5.3
/* 纯 JS，只需设置值 */
timeJs.set([t], 0)
/* 将数据从 CPU/JS 传递到 GPU/WGSL */
device.queue.writeBuffer(timeBuffer, 0, timeJs);</code>

パーティクルの位置 - WGSL ストレージ

粒子の位置を GPU からアクセス可能なメモリに直接保存して更新することで、GPU の大規模なマルチコア アーキテクチャを利用して粒子の位置を並行して更新できるようになります。

並列化は、コンピューティング シェーダーで宣言されたワーク グループ サイズを使用して調整されます。

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

@builtin(global_invocation_id) global_id : vec3 値はスレッド識別子を提供します。

定義により、global_invocation_id = workgroup_id * workgroup_size local_invocation_id - これは、パーティクル インデックスとして使用できることを意味します。

たとえば、10,000 個のパーティクルがあり、workgroup_size が 64 の場合、Math.ceil(10000/64) ワークグループをスケジュールする必要があります。コンピューティング パスが JS からトリガーされるたびに、次の量の作業を実行するように GPU に明示的に指示します:

<code>computePass.dispatchWorkgroups(Math.ceil(PARTICLE_COUNT / WORKGROUP_SIZE));</code>

PARTICLE_COUNT == 10000 および WORKGROUP_SIZE == 64 の場合、157 個のワークグループ (10000/64 = 156.25) が開始され、各ワークグループの local_invocation_id の計算範囲は 0 ～ 63 になります (workgroup_id の範囲は 0 ～ 157 です)。 ）。 157 * 64 = 1048 なので、最終的にはワークグループでもう少し多くの計算を行うことになります。冗長な呼び出しを破棄することでオーバーフローを処理します。

これらの要素を考慮したシェーダーの計算の最終結果は次のとおりです:

<code>@compute @workgroup_size(${WORKGROUP_SIZE})
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
  let index = global_id.x;
  // 由于工作组网格未对齐，因此丢弃额外的计算
  if (index >= arrayLength(&particles)) {
    return;
  }
  /* 将整数索引转换为浮点数，以便我们可以根据索引（和时间）计算位置更新 */
  let fi = f32(index);
  particles[index].position = vec2<f32>(
    /* 公式背后没有宏伟的意图 - 只不过是用时间+索引的例子 */
    cos(fi * 0.11) * 0.8 + sin((t + fi)/100)/10,
    sin(fi * 0.11) * 0.8 + cos((t + fi)/100)/10
  );
}
</f32></u32></code>

パーティクルはストレージ変数として定義されているため、これらの値は計算パス全体で保持されます。

頂点シェーダーのコンピュートシェーダーでパーティクルの位置を読み取ります

コンピューティング シェーダーのみがストレージに書き込むことができるため、コンピューティング シェーダーから頂点シェーダー内のパーティクルの位置を読み取るには、読み取り専用ビューが必要です。

以下は WGSL からの声明です:

<code>@group(0) @binding(0) var<uniform> t: f32;
@group(0) @binding(1) var<storage> particles : array<vec2>>;
/*
或等效：

@group(0) @binding(1) var<storage read=""> particles : array<vec2>>;
*/
</vec2></storage></vec2></storage></uniform></code>

コンピューティング シェーダーで同じ read_write スタイルを再利用しようとすると、エラーが発生します:

<code>var with 'storage' address space and 'read_write' access mode cannot be used by vertex pipeline stage</code>

頂点シェーダーのバインディング番号は、コンピューティング シェーダーのバインディング番号と一致する必要はないことに注意してください。頂点シェーダーのバインディング グループ宣言と一致する必要があるだけです。

<code>const renderBindGroup = device.createBindGroup({
  layout: pipeline.getBindGroupLayout(0),
  entries: [{
    binding: 0,
    resource: {
      buffer: timeBuffer
    }
  },
  {
    binding: 1,
    resource: {
      buffer: particleBuffer
    }
  }]
});</code>

GitHub サンプル コードで binding:2 を選択しました https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L70 - WebGPU によって課される制約の境界を調べるためだけに

シミュレーションを段階的に実行します

すべての設定が完了すると、更新ループとレンダリング ループが JS で調整されます。

<code>/* 从 t = 0 开始模拟 */
let t = 0
function frame() {
  /*
    为简单起见，使用恒定整数时间步 - 无论帧速率如何，都会一致渲染。
  */
  t += 1
  timeJs.set([t], 0)
  device.queue.writeBuffer(timeBuffer, 0, timeJs);

  // 计算传递以更新粒子位置
  const computePassEncoder = device.createCommandEncoder();
  const computePass = computePassEncoder.beginComputePass();
  computePass.setPipeline(computePipeline);
  computePass.setBindGroup(0, computeBindGroup);
  // 重要的是要调度正确数量的工作组以处理所有粒子
  computePass.dispatchWorkgroups(Math.ceil(PARTICLE_COUNT / WORKGROUP_SIZE));
  computePass.end();
  device.queue.submit([computePassEncoder.finish()]);

  // 渲染传递
  const commandEncoder = device.createCommandEncoder();
  const passEncoder = commandEncoder.beginRenderPass({
    colorAttachments: [{
      view: context.getCurrentTexture().createView(),
      clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 },
      loadOp: 'clear',
      storeOp: 'store',
    }]
  });
  passEncoder.setPipeline(pipeline);
  passEncoder.setBindGroup(0, renderBindGroup);
  passEncoder.draw(PARTICLE_COUNT);
  passEncoder.end();
  device.queue.submit([commandEncoder.finish()]);

  requestAnimationFrame(frame);
}
frame();</code>

結論

WebGPU は、ブラウザ内で大規模並列 GPU コンピューティングのパワーを解放します。

パスで実行されます。各パスには、メモリ バインディング (CPU メモリと GPU メモリのブリッジ) を備えたパイプラインを通じて有効化されたローカル変数があります。

コンピューティングの配信により、ワークグループを通じて並列ワークロードを調整できます。

多少の重いセットアップが必要ですが、ローカル バインディング/状態スタイルは WebGL のグローバル状態モデルに比べて大幅に改善されていると思います。使いやすくなり、同時に GPU コンピューティングのパワーを最終的に Web にもたらします。

以上がWebGPU チュートリアル: Web 上のコンピューティング、頂点、およびフラグメント シェーダーの詳細内容です。詳細については、PHP 中国語 Web サイトの他の関連記事を参照してください。