この記事は、WebGPUのごくごく基本的なことについて説明します。

WebGPUは、２つのことをやるためのAPIです。

1. 三角形/点/直線を、テクスチャに描く

2. GPU上で、計算を実行する

以上が、WebGPU APIでできること、です。

それ以外のことは、「WebGPU APIがやること」ではありません。あなたがやることです。 「WebGPUを学ぶ」というのは、「フレームワークの使い方を学ぶ」ようなことではなく、むしろJavaScriptやRust、C++のような「コンピュータ言語を学ぶ」ことに似ています。 基本を学んだその先の、基本を使って何を作るかは、全てあなたの創造力にゆだねられます。

WebGPUは極端に「低レベルなAPI」です。小さなサンプルプログラムを作る場合であっても、とてもたくさんのコードや、シビアなデータ構造の実装が必要となります。 たとえば、three.jsはWebGPUをサポートしていますが、そのサイズは600Kbytesにもなります。MinifyしたJavaSciprtコードの状態、もっと言えば、ローダや入力コントロール、ポストプロセスなどの機能を除いた、コア部分だけで、です。 TensorFlow.jsのWebGPUバックエンドでも同様です。こちらはMinifyした状態で500k程度となっています。

「画面に何か表示したい」といった観点で言えば、WebGPUを直接使うのではなく、ライブラリを使う方が、遥かに、良いです。

別の観点で、たとえば、「特定のユースケース向けのプログラムを作りたい」、「既存のライブラリに不満があって、改造したい」、「どういう仕組みか知りたい！全部知りたい！」、という場合は、WebGPUを直接扱うのが適している、と言えるでしょう。この記事はそんなあなたのためのものです。読み進めてください！

初めの一歩

「WebGPUの学習をどこから始めるか」は、案外難しい問題です。 WebGPUはある意味では、非常に単純な仕組みである、と言えます。 WebGPUは、「GPU上で３種類の関数を実行する」ことしかしません。「頂点シェーダ、フラグメントシェーダ、コンピュートシェーダ」の３種類です。

頂点シェーダ(Vertex Shader)は、頂点の計算をします。この関数は、頂点の位置を返します。得られた頂点の位置を３つ分使って、三角形が描画されます^[1]。

フラグメントシェーダ(Fragment Shader)は、色^[2]を計算します。GPUは、描画されるピクセル１個について１回、フラグメントシェーダを実行します。各フラグメントシェーダは、その１個のピクセルの色を返します。

コンピュートシェーダ(Compute Shader)は、もっと一般的な用途で利用されます。コンピュートシェーダは実質上、単なる関数であると考えてよいです。「これを何回実行します」という宣言を付けて実行されます。GPUはコンピュートシェーダを何度も呼び出すに際して、「何回目の呼び出しであるか」の情報をコンピュートシェーダに渡すので、その値を使って各呼び出しで挙動が変化するような書き方をする事ができます。

注意深い人は、これがJavaScriptの array.forEach や、 array.map といった仕組みに似ていることに気付くかもしれません。

GPU上で実行される関数は、あくまで関数にすぎません。JavaScriptの関数と同じように考えてもよいです。

違うのは、「GPU上で実行される」という部分です。このため、その関数から利用するデータは、あらかじめGPU上に、バッファやテクスチャの形にして置いておく必要があります。また、その関数が出力するデータも同様です。データは、あらかじめGPU上に確保したバッファやテクスチャに対して、書き出されます。

関数が使う「入出力データのありか」は、「バインディング」や「ロケーション」という形で、関数のコード中に記述します。 JavaScript側では、「バッファ」や「テクスチャ」といった「実際のデータ」と、「バインディング」や「ロケーション」といった「ありかの情報」を関連付けします。 これらの設定が終わったら、GPUに対して関数を実行する命令を発行します。

この話は、図にしてみると分かりやすいかも知れません。以下の図は、頂点シェーダとフラグメントシェーダを使って三角形を描くWebGPUアプリケーション、の実行に必要な設定を、単純化して表現したものです。

この図からは、次のようなことが読み取れます。

• Pipeline(パイプライン)：パイプラインは、頂点シェーダとフラグメントシェーダを持ちます。コンピュートシェーダを使う場合も、ここに設定します。

• 各シェーダが参照するリソース(buffers, textures, samplers)は、Bind Groups(バインドグループ)を介して、間接的に指定されます。

• パイプラインは、バッファを参照するアトリビュートを関連付けします。この関連付けは間接的なもので、実際には内部ステータスによって確定します。

• アトリビュートは、バッファからデータを引き出して、データは頂点シェーダで利用されます。

• 頂点シェーダは、フラグメントシェーダに対してデータを与える事ができます。

• フラグメントシェーダは、テクスチャに対して結果を書き出します。この書き出しは、render pass descriptionを通じて、間接的に行われます。

GPU上でシェーダを動かすためには、以上のような各種リソースを生成し、それぞれのステートを適切に設定する必要があります。「生成」は簡単な話です。注意すべき点があるとすれば、WebGPUでは一度生成したリソースは、ほとんどの場合、更新できない、というところでしょうか。リソースの中身を変更することはできますが、サイズ、用途、フォーマットといった項目については、変更ができません。これらの項目を変更したい場合は、リソースを新たに生成して、古い物を破棄(destroy)します。

コマンドバッファを生成します。いくつかのステートは、コマンドバッファ生成時に設定され、コマンドバッファが実行されます。コマンドバッファは、名前通り、コマンドのバッファです。

エンコーダを生成します。エンコーダは、コマンドを、コマンドバッファへエンコードします。以上を実行したら、エンコーダをfinishします。これによって、生成されたコマンドバッファが得られます。

コマンドバッファをsubmitします。これによって、WebGPUがコマンドを実行できます。

以下は疑似コードです。コマンドバッファをエンコードしてコマンドバッファが生成される様子を示しています。

encoder = device.createCommandEncoder()
// 何かを描画する
{
  pass = encoder.beginRenderPass(...)
  pass.setPipeline(...)
  pass.setVertexBuffer(0, …)
  pass.setVertexBuffer(1, …)
  pass.setIndexBuffer(...)
  pass.setBindGroup(0, …)
  pass.setBindGroup(1, …)
  pass.draw(...)
  pass.end()
}
// 何か別なものを描画する
{
  pass = encoder.beginRenderPass(...)
  pass.setPipeline(...)
  pass.setVertexBuffer(0, …)
  pass.setBindGroup(0, …)
  pass.draw(...)
  pass.end()
}
// 何やら計算する
{
  pass = encoder.beginComputePass(...)
  pass.beginComputePass(...)
  pass.setBindGroup(0, …)
  pass.setPipeline(...)
  pass.dispatchWorkgroups(...)
  pass.end();
}
commandBuffer = encoder.finish();

コマンドバッファの生成が済んだら、submit(送信)して実行できます。

device.queue.submit([commandBuffer]);

上の例のコマンドバッファでは、いくつかのdrawコマンドが登場しています。 各コマンドが実行されるとinternal state(内部ステート)が設定されていきます。 drawコマンドは、GPUに頂点シェーダを実行させます(間接的にフラグメントシェーダも実行されます)。dispatchWorkgroupコマンドは、GPUにコンピュートシェーダを実行させます。

このように、WebGPUでは「GPUにステートを設定、それを元にイメージを描く」という手順を踏みます。 ここまでの説明で、GPUでの描画に必要となるステートのイメージは、心に描けたでしょうか。

最初の方でも書きましたが、WebGPUは２つの基本的なことをやるAPIです。

1. 三角形/点/直線を、テクスチャに描く

2. GPU上で、計算を実行する

ここからは、この２つのことを実際にやる、小さなサンプルプログラムを作っていきます。別の記事では、プログラムにデータを与える様々な方法を紹介していくことになりますが、データを与える相手は結局のところ「この２つ」です。

改めて注意しますが、これからやるのは、ごくごく基本的なことです。 ここから上へ上へと積み上げていくために必要な、だいじな基礎部分です。 基礎部分が積み上がったら、GPUを使う典型的なアプリケーション、2Dグラフィクスや3Dグラフィクスのようなプログラムでその基礎がどう活用されるのか、紹介していきます。

三角形を、テクスチャに描く

WebGPUは、テクスチャに三角形を描くことができます。 この記事では、テクスチャは「二次元の長方形に並んだピクセル^[3]」としておきます。 HTMLの<canvas>要素ははWebページ中にテクスチャを提供するためのものです。WebGPUでは、canvasにテクスチャを要求して、そのテクスチャにレンダリングすることができます。

WebGPUで三角形を描くためには、２つの「シェーダ」を用意する必要があります。 既に説明したように、シェーダはGPU上で実行される関数です。シェーダは以下の二種類です。

1. 頂点シェーダ(Vertex Shader)

   頂点シェーダは、頂点の位置を計算する関数です。得られた頂点の位置は、三角形、線、点を描画するために使用されます。

2. フラグメントシェーダ(Fragment Shader)

   フラグメントシェーダは、色を計算する関数です。得られた色は、三角形、線、点を構成するピクセルの色として使用されます。色以外のデータを得るためにも使用されます。

ではWebGPUで三角形を描く、小さなプログラムを作っていきましょう。

まず、三角形が表示される、HTMLのcanvas要素が必要です。

<canvas></canvas>

JavaScriptを書くための<script>要素も必要です。

<canvas></canvas>
+<script type="module">

... javascriptのコードはここに書く ...

+</script>

以降のJavaScriptコードはすべて、上で書いたscriptタグの内側に書きます。

WebGPUは非同期APIです。非同期APIは、async functionを使って記述するのが簡単です。 まずアダプタを要求して、得られたアダプタからデバイスを取得します。

async function main() {
  const adapter = await navigator.gpu?.requestAdapter();
  const device = await adapter?.requestDevice();
  if (!device) {
    fail('WebGPU対応ブラウザが必要です');
    return;
  }
}
main();

上のコードは、見た通りで説明は不要かも知れません。 まず、?. オプショナルチェーン演算子を使ってアダプタをリクエストしています。 このため、navigator.gpuが存在しない場合、adapterはundefinedになります。 navigator.gpuが存在する場合は、requestAdapterを呼びます。requestAdapterの結果は非同期に返されるので、awaitで待つ必要があります。アダプタは、特定の１つのGPUを表すものです。デバイスによっては、複数のGPUを持っている場合もあります。

得られたアダプタから、デバイスを取得します。ここでも?.を使います。従って、アダプタがundefinedである場合は、デバイスもundefinedとなります。 この時点でdeviceが設定されていない場合は、古いブラウザを使っている、ということです。

次はcanvasを参照してwebgpuコンテキストを生成します。このコンテキストからは、シェーダの描画対象となるテクスチャが得られます。ブラウザはこのテクスチャを使って、Webページ上のcanvasに描画結果を表示します。

  // canvasからWebGPUコンテキストを取得して、configureする
  const canvas = document.querySelector('canvas');
  const context = canvas.getContext('webgpu');
  const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
  context.configure({
    device,
    format: presentationFormat,
  });

このコードも、見た通りで説明は不要かも知れません。 まず、canvasから"webgpu"コンテキストを取得しています。 推奨のcanvasフォーマット(preferred canvas format)をシステムに問い合わせして、取得しています。 これは"rgba8unorm"、"bgra8unorm"のいずれかです。これが何かはひとまず置くとして、 この問い合わせをすることで、そのシステムにおいて最速な処理方法が選択できます。 WebGPUコンテキストに対して、configureを使ってこのformatを渡しています。 同時にdeviceを渡しています。これでさっき取得したデバイスとWebGPUコンテキストが関連付けられます。

次は、シェーダモジュールを生成します。 シェーダモジュールは、１つ以上のシェーダを持つコンテナです。 今回は、頂点シェーダを１つ、フラグメントシェーダを１つ、都合２つのシェーダをシェーダモジュールに持たせることにします。

  const module = device.createShaderModule({
    label: 'our hardcoded red triangle shaders',
    code: `
      @vertex fn vs(
        @builtin(vertex_index) vertexIndex : u32
      ) -> @builtin(position) vec4f {
        let pos = array(
          vec2f( 0.0,  0.5),  // top center
          vec2f(-0.5, -0.5),  // bottom left
          vec2f( 0.5, -0.5)   // bottom right
        );

        return vec4f(pos[vertexIndex], 0.0, 1.0);
      }

      @fragment fn fs() -> @location(0) vec4f {
        return vec4f(1.0, 0.0, 0.0, 1.0);
      }
    `,
  });

シェーダはWGSL (WebGPU Shading Language)と呼ばれる言語で記述されます。WGSLは、しばしば「ウィグシル」と発音されているようです。 WGSLは「強い型付け言語」の一種です。詳細については別途「WGSLについて」で説明することとして、ここでは、大ざっぱに見通せる程度のそれなりな説明に留めたいと思います。

上のソースコードでは、vsと名付けた関数を、@vertex属性を付けて宣言しています。 @vertex属性は、この関数が頂点シェーダであることを示します。

      @vertex fn vs(
        @builtin(vertex_index) vertexIndex : u32
      ) -> @builtin(position) vec4f {
         ...

この関数vsでは、vertexIndexと名付けた引数を１つ、受け付けています。 vertexIndexはu32、つまり32ビット符号なし整数型の変数です。 vertexIndexの値は、ビルトインであるvertex_indexから取得されます。 vertex_indexビルトインは、繰り返しのループカウンタ(iteration number)のようなものです。 JavaScriptで言えばArray.map(function(value, index) { ... })のindexに相当します。 GPUにdrawコマンドを渡す際に「10回繰り返す」ように指定した場合、 シェーダ実行１回目では、vertex_indexの値は0、 シェーダ実行２回目では、vertex_indexの値は1 ……などとなります^[4]。

関数vsは、vec4f型の値を返り値として持つように宣言されています。 vec4f型は、32ビット浮動小数点数の値を4つ持つベクトルです。これは、「４つの値を持つ配列」とか、{x: 0, y: 0, z: 0, w: 0}のような「４つのプロパティを持つオブジェクト」のようなものです。 この返り値は、positionビルトインに代入されます。 今回使用する"triangle-list"モードの場合、頂点シェーダが３回呼ばれて３回分のposition が得られるたびに、三角形が描画されます。

WebGPUでは、positionはclip space(クリップ空間)の座標として扱われます。クリップ空間は、 Xの左端が-1.0、右端が+1.0、Yの下端が-1.0、上端が+1.0です。 これは描画対象のテクスチャのサイズと無関係で、常に-1.0から1.0の範囲です。

関数vsの中では、「vec2f型の値３つを持つ配列(array)」が宣言されています。 vec2fは、32ビット浮動小数点数の値２つを持ちます。

        let pos = array(
          vec2f( 0.0,  0.5),  // top center
          vec2f(-0.5, -0.5),  // bottom left
          vec2f( 0.5, -0.5)   // bottom right
        );

関数vsの最後の部分では、vertexIndexを使って、配列中の３つの値を、呼び出し元に返しています。 関数vsの返り値の型はvec4fとしたので、返すべき情報は浮動小数点数４つ、です。 一方で、先ほど宣言したposはvec2fの配列、その配列要素であるpos[vertexIndex]はvec2f、 つまり浮動小数点数２つ、です。このため、コードでは0.0と1.0を付け足して、浮動小数点数４つ、 として数を合わせています。

        return vec4f(pos[vertexIndex], 0.0, 1.0);

シェーダモジュールでは、関数vsのほかに、fsと名付けた関数の宣言も行なっています。 関数fsは、@fragment属性が付けられて宣言されているため、フラグメントシェーダとして扱われます。

      @fragment fn fs() -> @location(0) vec4f {

関数fsは、引数なし、返り値はvec4f型でlocation(0)に返されます。 location(0)というのは、「一つめのレンダーターゲット」という意味で、結果はそこに書き込まれます。 今回は「一つめのレンダーターゲット」を、「用意しておいたcanvasのテクスチャ」としたいのですが、その設定の仕方については後ほど説明します。

        return vec4f(1, 0, 0, 1);

このコードでは1, 0, 0, 1を返しています。これは「赤」です。 WebGPUでの色指定は、一般には0.0から1.0の範囲値を４つ使います。 ４つの値はそれぞれ「赤、緑、青、アルファ(透明度)」に対応します。

フラグメントシェーダは、GPUが三角形をラスタライズ(ピクセルの集まりとして描画)する際に呼ばれます。 １つのピクセルに対してフラグメントシェーダが１回呼ばれて、そのピクセルの色が決まります。 今回の例では、三角形を構成する全てのピクセルで「赤」が返されます。

WebGPUのコードを書く上で、私が重要と考えることについて触れておきます。 「label(ラベル)」についてです。 WebGPUで生成するオブジェクトは、ほぼすべてについてlabelを付けることができます。 ラベルを付けることはまったく必須ではない、のですが、 生成するもの全てにラベルを付けることはbest practice(ベストプラクティス。よい習慣)であると言えます。 エラーが発生した場合、WebGPU実装はエラーメッセージを出力します。 この時、エラーメッセージにはエラー箇所に関連したラベルが記述されます。

標準的なWebGPUアプリケーションでは、バッファ、テクスチャ、シェーダモジュール、パイプラインといったオブジェクトが100個、1000個と登場します。 もし、プログラミングをしていて発生したエラーメッセージが "WGSL syntax error in shaderModule at line 10" といったものだったらどうでしょう？ シェーダモジュールが100個あったら、どのシェーダモジュールか特定できるでしょうか？ 一方で、これが、 "WGSL syntax error in shaderModule('our hardcoded red triangle shaders') at line 10 といったものだったらどうでしょう？ 問題判別に費やす膨大な時間が、節約できるのではないでしょうか？

さて。ここまででシェーダモジュールが生成できました。 次にすべきことは、レンダーパイプラインを作ることです。

  const pipeline = device.createRenderPipeline({
    label: 'our hardcoded red triangle pipeline',
    layout: 'auto',
    vertex: {
      module,
      entryPoint: 'vs',
    },
    fragment: {
      module,
      entryPoint: 'fs',
      targets: [{ format: presentationFormat }],
    },
  });

この部分については、今回のサンプルプログラムでは見るべきところはないかも知れません。 layoutを'auto'に設定しています。これはデータのレイアウトを、シェーダのコードの内容からWebGPUが自動で設定する、という意味です。ただ、今回はレイアウトすべきデータ自体がありません。

次の部分では、シェーダモジュールに記述された関数vsを頂点シェーダとして使う、関数fsをフラグメントシェーダとして使う、ということをレンダーパイプラインに教えています。 同時に、１つめのレンダーターゲットにフォーマットを指定しています。 「レンダーターゲット」というのは、書き込み先とするテクスチャのことです。 パイプラインを生成し、このパイプラインの書き出し先となるテクスチャのフォーマットを指定する必要があります。

配列targetsの最初の(0番目の)要素は、先にフラグメントシェーダの返り値の設定で記述した「location(0)」に当たるものです。canvasとの関連付けはもう少し後で行ないます。

次はGPURenderPassDescriptorを設定します。 これは、描画先とするテクスチャの指定と、そのテクスチャをどう扱うかの設定です。

  const renderPassDescriptor = {
    label: 'our basic canvas renderPass',
    colorAttachments: [
      {
        // view: <- viewプロパティの設定はレンダリングするタイミングで行なう。
		clearValue: [0.3, 0.3, 0.3, 1],
        loadOp: 'clear',
        storeOp: 'store',
      },
    ],
  };  

GPURenderPassDescriptorのプロパティcolorAttachmentsは配列です。 この配列は、「描画対象となるテクスチャ」です。 また、「各テクスチャをどう扱うか」という情報も記述されます。 このタイミングでは「描画対象となるテクスチャ」の指定は保留して、「各テクスチャをどう扱うか」の設定を行っています。 clearValueは、背景色、つまり全体を一色で塗りつぶす際の指定色です。暗めの灰色[0.3, 0.3, 0.3, 1]を指定しています。 loadOp: 'clear'は、「描画開始前にテクスチャ全体を背景色でクリアする」という設定です。loadOp: 'load'とした場合は、「その時点のテクスチャの内容をGPUにロードして、そこに上書きで描画していく」という意味になります。 storeOp: 'store'は、描画内容をテクスチャに保存する、という意味です。storeOp: 'discard'とすると、描画内容を破棄します。'discard'がどういう場面で役に立つのか、については別の記事、「マルチサンプリング」で説明します.

さて。レンダリングする(描画する)時が来ました。

  function render() {
    // canvasのコンテキストから、カレントテクスチャを得る。
	// それをレンダーパスに設定して、描画対象として指定する。
    renderPassDescriptor.colorAttachments[0].view =
        context.getCurrentTexture().createView();

    // コマンドエンコーダを生成する。コマンドのエンコードができる状態にする。
	const encoder = device.createCommandEncoder({ label: 'our encoder' });

    // レンダーパスのエンコーダを生成する。そこへコマンドを並べて、描画手順をエンコードする。
    const pass = encoder.beginRenderPass(renderPassDescriptor);
    pass.setPipeline(pipeline);
    pass.draw(3);  // 頂点シェーダを３回呼び出す
    pass.end();

    const commandBuffer = encoder.finish();
    device.queue.submit([commandBuffer]);
  }

  render();

最初にcontext.getCurrentTexture()を呼んで、canvasが持っているテクスチャを取得しています。createViewでは、テクスチャの一部の範囲だけを切り出す指定ができますが、ここでは引数なし＝デフォルトの範囲、としています。 今回は、配列colorAttachmentsの要素はひとつだけとしていました。このcolorAttachments[0]に、先ほど取得したcanvasのテクスチャ(texture view)を設定します。 先だって述べたように、このcolorAttachments[0]は、フラグメントシェーダの返り値の設定で記述したlocation(0)に対応するものです。

次の部分では、コマンドエンコーダを用意しています。コマンドエンコーダは、コマンドバッファを生成するために使われます。コマンドエンコーダを使って、各種コマンドをコマンドバッファに並べて行きます。"submit"でコマンドバッファを送信すると、コマンドが実行されます。

次は、コマンドエンコーダのbeginRenderPassを使って、レンダーパスエンコーダを生成しています。レンダーパスエンコーダは、レンダリング関連のコマンドを生成することに特化したエンコーダです。 レンダーパスエンコーダにrenderPassDescriptorを渡すことで、描画対象とするテクスチャを指定しています。

setPipelineコマンドでパイプラインをセット、次のdraw(3)コマンドで、頂点シェーダを３回実行します。デフォルトでは、頂点シェーダが３回実行されるたびに、３つの頂点シェーダが返した３つの点を結ぶ三角形が描画されます。

レンダーパスを終了し、エンコーダをfinishしています。finishを実行することで、先ほどコマンドを並べて定義した手順が入ったコマンドバッファが得られます。最後に、このコマンドバッファをsubmit(送信)し、レンダリングを実行しています。

drawが実行される段階での各種ステートは以下のようになっています。

今回の例では、入力としては、テクスチャ、バッファ、バインドグループは一切ありません。使っているのは、パイプライン、頂点シェーダ、フラグメントシェーダ、出力先のcanvasテクスチャを指定しているレンダーパスディスクリプタ、です。

ここまで書いてきたコードの実行結果は以下のようになります。

今回使ったsetPipelineやdrawといったAPIは「コマンドバッファにコマンドを追記するだけのもの」である、という点は、重要です。これらのAPIを実行した時点では、コマンドの実行は起きません。コマンドの実行は、デバイスキューにコマンドバッファをsubmit(送信)して初めて行なわれます。

WebGPUは、頂点シェーダを３回呼び出して得た頂点を３つを集めて、三角形をラスタライズします。 WebGPUはこの際、「各ピクセルの中心点が三角形の内側にあるかどうか」を見ることで、どのピクセルを描画するか判断しています。 描画することになったピクセルには、フラグメントシェーダを呼んで得られた色が塗られることになります。

例えば、15x11ピクセルのテクスチャに三角形を描画する場合、ラスタライズは以下図のような仕組みで行われます。

ここまで、WebGPUのごく小さなサンプルプログラムを作ってきました。お気づきと思いますが、シェーダコード中に三角形がハードコードされた今回のコードは、柔軟性が全くありません。一般にはシェーダの外からデータを与えることになるわけですが、それについては今後の記事で紹介することとします。

今回のサンプルプログラムで示したかったことをまとめます。

• WebGPUはシェーダを実行するのが仕事。それ以外は、全てあなたの仕事
• シェーダはシェーダモジュールで指定する。シェーダモジュールはパイプラインの中に置く
• WebGPUは三角形を描くことができる
• WebGPUはテクスチャを描画できる(今回はテクスチャをcanvasから取得した)
• WebGPUは、エンコードしたコマンドをsubmitすると動く

GPU上で、計算を実行する

GPU上で計算を実行する、簡単なサンプルプログラムを作っていきます。

WebGPUデバイスの取得の手順は先ほどと同じです。

async function main() {
  const adapter = await navigator.gpu?.requestAdapter();
  const device = await adapter?.requestDevice();
  if (!device) {
    fail('need a browser that supports WebGPU');
    return;
  }

違ってくるのはシェーダモジュールの部分からです。

  const module = device.createShaderModule({
    label: 'doubling compute module',
    code: `
      @group(0) @binding(0) var<storage, read_write> data: array<f32>;

      @compute @workgroup_size(1) fn computeSomething(
        @builtin(global_invocation_id) id: vec3u
      ) {
        let i = id.x;
        data[i] = data[i] * 2.0;
      }
    `,
  });

最初の部分で、storageタイプの変数dataを宣言しています。読み出し可能、書き込み可能としています。

      @group(0) @binding(0) var<storage, read_write> data: array<f32>;

変数の型はarray<f32>としています。これは、32ビット浮動小数点数の配列です。 また、この変数をbinding(0)、@group(0)と関連付けています。それぞれ、バインドロケーションの0番、バインドグループの0番です。

次に、関数computeSomethingを宣言しています。宣言では@compute属性を付けることで、この関数をコンピュートシェーダとしています。

      @compute @workgroup_size(1) fn computeSomething(
        @builtin(global_invocation_id) id: vec3u
      ) {
        ...

コンピュートシェーダでは、「ワークグループサイズ」を宣言する必要があります。「ワークグループサイズ」については後述します。ここではとりあえず属性として@workgroup_size(1)を設定、つまりワークグループサイズを1としています。

関数computeSomethingの宣言では、vec3u型の引数idを設定しています。 vec3uは、符号なし32ビット整数値を３持つベクトルです。 このidは、頂点シェーダの宣言の引数と同様、繰り返し番号です。違っているのは、このコンピュートシェーダの引数はvec3uなので三次元である(つまり、値を３つ持つ)、ということです。idの値は、global_invocation_idビルトインから取得するように宣言しています。

コンピュートシェーダの引数の扱いについては、かなり大雑把ですが、以下のような感じの仕組みだと考えてみると良いと思います。

// 疑似コード！
function dispatchWorkgroups(width, height, depth) {
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const workgroup_id = {x, y, z};
        dispatchWorkgroup(workgroup_id)
      }
    }
  }
}

function dispatchWorkgroup(workgroup_id) {
  // from @workgroup_size in WGSL
  const workgroup_size = shaderCode.workgroup_size;
  const {x: width, y: height, z: depth} = workgroup_size;
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const local_invocation_id = {x, y, z};
        const global_invocation_id =
            workgroup_id * workgroup_size + local_invocation_id;
        computeShader(global_invocation_id)
      }
    }
  }
}

今回は@workgroup_size(1)としたので、実質的には以下のように考えても同じです。

// 疑似コード！
function dispatchWorkgroups(width, height, depth) {
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const workgroup_id = {x, y, z};
        dispatchWorkgroup(workgroup_id)
      }
    }
  }
}

function dispatchWorkgroup(workgroup_id) {
  const global_invocation_id = workgroup_id;
  computeShader(global_invocation_id)
}

最後に、idのプロパティxを配列dataの添え字として、その値(つまりdata[id.x])を２倍しています。

        let i = id.x;
        data[i] = data[i] * 2.0;

この例ではiは、繰り返し番号の３つの値の最初の１つだけ使っています。

以上で、パイプラインに必要となるシェーダが作成できました。

  const pipeline = device.createComputePipeline({
    label: 'doubling compute pipeline',
    layout: 'auto',
    compute: {
      module,
      entryPoint: 'computeSomething',
    },
  });

このパイプラインでは、computeステージについて先ほど作ったシェーダモジュールmoduleを使う、起動時にはcomputeSomething関数を呼ぶ、といったことを言っています。 layoutは今回も'auto'としておきます。データのレイアウトをシェーダのコードからWebGPUが自動で判別して設定します。^[5]

次は、データを用意します。

  const input = new Float32Array([1, 3, 5]);

今回は、データはJavaScript側にのみ記述しています。 WebGPUでこのデータを利用するには、GPU上にバッファを用意して、そのバッファにデータをコピーする、必要があります。

  // 計算の入出力に使うバッファを、GPU上に用意する。
  const workBuffer = device.createBuffer({
    label: 'work buffer',
    size: input.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
  });
  // JavaScript側で用意した入力データを、GPU上のバッファへコピーする。
  device.queue.writeBuffer(workBuffer, 0, input);

上のコードでは、まずdevice.createBufferでバッファを生成しています。 sizeはバッファのサイズで、単位はバイトです。 今回の例ではFloat32Arrayの値を３つ扱うので12バイトになります。 Float32Arrayについては「メモリレイアウト」で説明しているので、Float32Arrayなじみがないという人は読んでみると良いでしょう。

WebGPUで利用するバッファでは、必ずusageを指定する必要があります。 usageに指定できるフラグは多数あり、同時に指定できないフラグなどもあります。 今回のケースでは、

• GPUBufferUsage.STORAGEを指定して、storageとして利用できるようにしています。この設定は、シェーダの方でvar<storage,...>と設定したものと対応しています。
• GPUBufferUsage.COPY_DSTを指定して、このバッファをデータのコピー先とできるようにしています。
• GPUBufferUsage.COPY_SRCを指定して、このバッファをデータのコピー元とできるようにしています。

ここでひとつ注意なのですが、WebGPUのバッファの中身は、JavaScript側から直接見ることができません。 直接見るのではなく、WebGPUのバッファ上のデータを、JavaScript側へマップする必要があります。というのも、そうしないと、読みだしている最中にバッファの内容が更新中であった場合に問題が起きるからです。また、そもそもバッファのデータはCPU側のメモリには存在せず、GPU側のメモリにあるから、というのも理由の一つです。

WebGPUのバッファをJavaScript側にマップすることはできますが、マップできるようにしたバッファはそれ以外の操作はできません。 たとえば、今回作ったバッファはマップすることができません。usageをSTORAGEとしたためです。この設定のバッファをマップしようとするとエラーになります。

従って、JavaScript側から計算結果を見るためには、usageがSTORAGEではない別のバッファが必要になります。計算が完了した後に、「計算結果を書き込むためのバッファ」から、「JavaScript側へマップ可能なバッファ」へコピーします。

  // GPUの外から見えるように、計算結果をコピーする新たなバッファを、GPU上に用意する
  const resultBuffer = device.createBuffer({
    label: 'result buffer',
    size: input.byteLength,
    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
  });

MAP_READは、データをGPUの外へマップできるようにする、という意味です。

これらのバッファの情報をシェーダに伝えるため、バインドグループ(bindGroup)を作ります。

  // 計算をする際にどのバッファを使えばよいかシェーダに指示するため、
  // bindGroupを設定する。
  const bindGroup = device.createBindGroup({
    label: 'bindGroup for work buffer',
    layout: pipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: { buffer: workBuffer } },
    ],
  });

pipeline.getBindGroupLayoutで、パイプラインからbindGroupを取得しています。 bindGroupのエントリを設定します。 pipeline.getBindGroupLayout(0)の0は、シェーダで記述した@group(0) に相当します。 entriesの{binding: 0 ...は、シェーダで記述した@group(0) @binding(0) に相当します。

次はコマンドのエンコードです。

  // 計算用のコマンドをエンコードする
  const encoder = device.createCommandEncoder({
    label: 'doubling encoder',
  });
  const pass = encoder.beginComputePass({
    label: 'doubling compute pass',
  });
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(input.length);
  pass.end();

createCommandEncoderでコマンドエンコーダを生成、 beginComputePassでコンピュートパスを開始、 setPipeline(pipeline)でパイプラインをセット、 setBindGroup(0, bindGroup)でbindGroupをセットしています。この0は、シェーダで記述した@group(0)に対応するものです。 dispatchWorkgroupsを呼んでいます。input.lengthの値は3なので、これはWebGPUに対して「コンピュートシェーダを３回呼べ」という命令になります。 endでコンピュートパスを終了しています。

これはdispatchWorkgroupsの実行時点の、各種ステートです。

計算が終わったら、workBufferに書き込まれたデータをGPUの外へ持ち出すためにresultBufferへコピーします。

  // 「得られた結果をマップ可能なバッファへコピーするコマンド」をエンコードする。
  encoder.copyBufferToBuffer(workBuffer, 0, resultBuffer, 0, resultBuffer.size);

エンコーダをfinishして、コマンドバッファを取得します。 得られたコマンドバッファを、GPUに対してsubmitします。

  // コマンドのエンコードを完了。コマンドバッファをGPUへsubmitする。
  const commandBuffer = encoder.finish();
  device.queue.submit([commandBuffer]);

resultBufferをマップして、データのコピーを取得します。

  // 計算結果を読み出す。
  await resultBuffer.mapAsync(GPUMapMode.READ);
  const result = new Float32Array(resultBuffer.getMappedRange());

  console.log('input', input);
  console.log('result', result);

  resultBuffer.unmap();

バッファをマップするためにmapAsyncを使用します。この関数を利用する際は、処理の完了をawaitで待つ必要があります。 マップの処理が完了したらresultBuffer.getMappedRange()でArrayBufferを取得しています。引数なしで読んだ場合、ArrayBuffer全体が返されます。 得られたバイト列であるArrayBufferを、型付き配列であるFloat32Arrayのビューとすることで、シェーダの計算結果を数値として見ることができます。 一つ重要な注意点なのですが、getMappedRangeが返したArrayBufferが有効なのは、unmapを呼ぶまでの間だけ、となっています。unmapを実行すると、ArrayBufferのlengthは0となり、データにアクセスすることはできなくなります。

これが実行結果です。３つの数値をGPUに渡して、GPU上のシェーダで計算して、二倍になった値をJavaScript側へ取得して、表示しています。

コンピュートシェーダの本格的な使い方は別の記事で紹介しますが、ここまでで、WebGPUが何をするものであるか、一通りのことが伝わったのではないかと思います。この先は、全てあなた次第です。 WebGPUの学習はプログラム言語の学習のようなものです。WebGPUはいくつかの基本機能を提供するもので、それ以上の部分はあなたの創造性にゆだねられています。

WebGPUのプログラミングを特別なものにしているのは、頂点シェーダ、フラグメントシェーダ、コンピュートシェーダで、これらのシェーダがGPU上で動く、というところです。 GPUは、ものによっては１つのGPUで10000個にも及ぶプロセッサを内包しています。 これは、ポテンシャルとしては10000回の計算を並列で実行できるということです。並列性の規模は、CPUとは実に三桁もの違いがあります。

簡単なcanvasリサイズ

別の話題に移る前に、三角形を描くサンプルを利用して「canvasをリサイズする仕組み」を紹介しておきます。canvasサイズの変更については「canvasのリサイズ」で記事一本丸ごと使って詳しく説明しています。 ここで紹介するのは簡易版です。

最初に、CSSを追加してcanvasがページ全体を覆うようにします。

<style>
html, body {
  margin: 0;       /* デフォルトのmarginをなくす          */
  height: 100%;    /* html,bodyをページ全体に合わせる     */
}
canvas {
  display: block;  /* canvasの挙動をブロック要素扱いにする */
  width: 100%;     /* canvasを自身のコンテナに合わせる     */
  height: 100%;
}
</style>

このCSSは、canvasをページ全体に表示するものです。表示サイズを設定するもので、canvas自体の解像度は変化しません。 このため、別ウィンドウで表示して最大化すると、三角形が何やらギザギザになるのが分かると思います。

<canvas>タグの解像度は、デフォルトでは300x150ということになっています。 これを、表示サイズに合わせてcanvasの解像度が変化するようにします。

これにはいくつかやり方がありますが、ここではResizeObserverを使ってみます。 まず、ResizeObserverクラスのオブジェクトを生成。監視対象がリサイズされるたびに処理を行うメソッドを追加。監視対象のHTML要素を指定、しています。

    ...
-    render();

+    const observer = new ResizeObserver(entries => {
+      for (const entry of entries) {
+        const canvas = entry.target;
+        const width = entry.contentBoxSize[0].inlineSize;
+        const height = entry.contentBoxSize[0].blockSize;
+        canvas.width = Math.max(1, Math.min(width, device.limits.maxTextureDimension2D));
+        canvas.height = Math.max(1, Math.min(height, device.limits.maxTextureDimension2D));
+        // re-render
+        render();
+      }
+    });
+    observer.observe(canvas);

上のコードではentriesの全ての要素をforで走査していますが、実際に操作対象となるのはobserveで指定したcanvas一つだけです。canvasのサイズには、デバイスによる上限の制限があるので、それに配慮しています。この上限を超えてしまうと、WebGPUは、大きすぎるテクスチャを生成しようとしてエラーを出力することになります。また、サイズが0の場合も同様にエラーとなります。 これについては、「canvasのリサイズ」の記事で説明します。

canvasの解像度が適切に更新できたところでrenderを呼んで再描画しています。元のコードでは無条件でrenderを呼んでいましたが、その部分は不要なのでコードから削除しています。ResizeObserverは、observeを開始した時点でのcanvas要素のサイズを報告するため、コールバックを必ず１回は実行するためです。

リサイズ後のサイズの新たなテクスチャの生成は、renderの中に書いたcontext.getCurrentTexture()が行なうので、observerは関与しません。

以下の記事では、シェーダにデータを渡す様々な方法、について説明しています。

• inter-stage variables
• uniforms
• storage buffers
• vertex buffers
• textures
• constants

また「WGSLの基礎」では、WGSLについて説明します。

上に記事順は、簡単なものから複雑なもの、の順です。

inter-stage variable(ステージ間変数)の利用は、WGSLのコードの中で完結できるので、 外部での設定の説明は必要がありません。WGSLのコードの変更だけで利用できます。 uniform(uniform変数)は実質上、グローバル変数です。三種類あるシェーダ(頂点シェーダ、フラグメントシェーダ、コンピュートシェーダ)の、いずれでも利用されます。 strage buffer(ストレージバッファ)の利用は、uniform bufferからstrage bufferへの移行であればとても簡単です。strage buffer記事冒頭で説明しています。 vertex buffer(頂点バッファ)は、頂点シェーダでのみ利用されます。これはそれなりに複雑です。利用するためにはWebGPUのデータレイアウトの記述が必要になります。 texture(テクスチャ)はさらに複雑です。利用するために設定が必要なオプションが山のようにあります。

これらの話題は退屈なので、読み飛ばしてしまうのも良いかも知れません。 WebGPUの学習の入口で飽きてしまうよりは良いです。 これらの基礎知識が必要である、と実感する場面まで進んでから、改めて読み直すなり、初めて読むなりするのが良いと思います。 基本的なことを押さえることができたら、実用的なテクニックに触れていきます。

もう一つ付け加えておきます。このWebサイトのプログラム例は全て、その場で編集、実行することができるようにしてあります。また、jsfiddle、codepen、あとstackoverflowへ持っていくことも"Export"ワンタッチでできます。

async function start() {
  if (!navigator.gpu) {
    fail('this browser does not support WebGPU');
    return;
  }

  const adapter = await navigator.gpu.requestAdapter();
  if (!adapter) {
    fail('this browser supports webgpu but it appears disabled');
    return;
  }

  const device = await adapter?.requestDevice();
  device.lost.then((info) => {
    console.error(`WebGPU device was lost: ${info.message}`);

    // 'reason' will be 'destroyed' if we intentionally destroy the device.
    if (info.reason !== 'destroyed') {
      // try again
      start();
    }
  });
  
  main(device);
}
start();

function main(device) {
  ... do webgpu ...
}