WebGL / WebGPU公開 約15分

WebGPU uniformの配列配置:vec3の12バイトと16バイトを分ける

三成分を二組送ったのに、二組目だけ違う値になる場合に。WGSLのvec3配列とJavaScriptのTypedArrayを対応させ、余白を省いた値のずれと、領域不足による接続失敗を分けて確認します。

この記事の目次
  1. 三成分の大きさと、次の要素までの間隔
  2. 同じ六つの数値を、二通りに並べる
  3. 領域が足りても、値の場所は間違えられる
  4. 六成分分だけ確保すると、接続時に止まる
  5. この16バイトを、すべての型へ当てはめない

二つの色をGPUへ渡したのに、二つ目だけ違う値になる。JavaScriptの配列を見ても六つの成分はそろっている。そんなときは、数値の順番だけでなく、次の要素が始まる位置を調べます。GPUは、JavaScript側の配列の区切りを自動で補ってはくれません。

対象はJavaScriptの非同期関数とTypedArrayを読める人です。WebGPUとWGSLを使い、架空の配色見本から六つの数値を読み戻します。描画する前に配置だけを調べる実験なので、canvasやアニメーションはありません。追加ライブラリも使いません。

三成分の大きさと、次の要素までの間隔

uniformは、シェーダーへ読み取り用の値を渡すための領域です。今回の型は、三つの32ビット浮動小数点数を持つvec3<f32>の二要素配列です。三成分そのものは12バイトですが、要素の開始位置には16バイト単位の整列が必要です。

要素の先頭から次の要素の先頭までの間隔をstrideと呼びます。この型では12を16の倍数へ切り上げ、strideは16バイトになります。WGSLの配置規則で、サイズと整列、配列のstrideが別の量として定義されています。

用途 先頭からのバイト位置 Float32Arrayの添字
一つ目の三成分 0、4、8 0、1、2
次の要素までの余白 12 3
二つ目の三成分 16、20、24 4、5、6
配列末尾の余白 28 7

必要な領域は32バイト、Float32Arrayなら8要素です。三成分だから毎回4番目に値を入れる、という意味ではありません。余白は今回読み取る値に含めず、二つ目の先頭を添字4へそろえるために残します。

三成分を詰めると二つ目が12バイトから始まる。WGSLの配列に合わせるには余白を挟み16バイトから始める

同じ六つの数値を、二通りに並べる

一つ目の色を0.25、0.5、1、二つ目を0.75、0.25、0.5とします。ここでは色を画面で比べず、GPUが読んだ成分をそのまま結果用バッファへ書き出します。色空間やディスプレイの違いが観察へ混ざらないようにするためです。

次をuniform.htmlへ保存します。フォルダ内でpython3 -m http.server 8768 --bind 127.0.0.1を実行し、http://127.0.0.1:8768/uniform.htmlを開きます。Python 3はローカル配信用で、終了はCtrl+Cです。WebGPUを利用できるブラウザーと、HTTPSまたは信頼されるローカル環境が必要です。

adapter取得の説明のとおり、APIがあっても利用可能なadapterを得られない場合があります。その場合は数値の結果を作らず、利用できない理由を文章で表示します。

<!doctype html>
<html lang="ja">
<head>
<meta charset="utf-8">
<meta name="viewport" content="width=device-width, initial-scale=1">
<title>二つの色の数値を読み戻す</title>
<style>
body { margin: 16px; font: 16px/1.8 sans-serif; }
main { max-width: 720px; margin: auto; }
button, select { font: inherit; margin-block: 8px; max-width: 100%; }
pre { white-space: pre-wrap; overflow-wrap: anywhere; }
:focus-visible { outline: 3px solid #245ba7; outline-offset: 3px; }
</style>
</head><body><main>
<h1>二つの色の数値を読み戻す</h1>
<p>色の見た目ではなく、GPUが読んだ六つの成分を調べます。</p>
<label for="mode">試す条件</label>
<select id="mode">
<option value="ok">16バイトごとに配置</option>
<option value="packed">隙間なく詰める</option>
<option value="short">領域を24バイトにする</option>
</select>
<button id="run" type="button">実行する</button>
<pre id="result" role="status">条件を選んで実行してください。JavaScriptが必要です。</pre>
</main><script type="module">
const button = document.querySelector('#run');
const result = document.querySelector('#result');
const mode = document.querySelector('#mode');
async function checked(device, label, work) {
  device.pushErrorScope('validation');
  let value;
  try { value = work(); }
  finally {
    const error = await device.popErrorScope();
    if (error) throw new Error(label + ': ' + error.message);
  }
  return value;
}
async function run(kind) {
  let device;
  let finished = false;
  let lost = false;
  const buffers = [];
  try {
    if (!navigator.gpu) throw new Error('WebGPUを利用できません。');
    const adapter = await navigator.gpu.requestAdapter();
    if (!adapter) throw new Error('利用できるadapterがありません。');
    device = await adapter.requestDevice();
    device.lost.then(() => {
      if (!finished) {
        lost = true;
        result.textContent = 'デバイスを失いました。再実行してください。';
      }
    });
    function buffer(size, usage) {
      const value = device.createBuffer({ size, usage });
      buffers.push(value);
      return value;
    }
    const module = device.createShaderModule({ code: `
struct Palette { colors: array<vec3<f32>, 2>, }
@group(0) @binding(0) var<uniform> palette: Palette;
@group(0) @binding(1) var<storage, read_write> result: array<f32, 6>;
@compute @workgroup_size(1) fn main() {
  for (var i = 0u; i < 2u; i++) {
    for (var c = 0u; c < 3u; c++) {
      result[i * 3u + c] = palette.colors[i][c];
    }
  }
}` });
    const pipeline = await device.createComputePipelineAsync({
      layout: 'auto', compute: { module, entryPoint: 'main' }
    });
    const data = new Float32Array(kind === 'short' ? 6 : 8);
    data.set([0.25, 0.5, 1], 0);
    data.set([0.75, 0.25, 0.5], kind === 'packed' || kind === 'short' ? 3 : 4);
    const input = buffer(data.byteLength, GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST);
    const output = buffer(24, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
    const readback = buffer(24, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
    device.queue.writeBuffer(input, 0, data);
    const group = await checked(device, '接続', () => device.createBindGroup({
      layout: pipeline.getBindGroupLayout(0),
      entries: [
        { binding: 0, resource: { buffer: input } },
        { binding: 1, resource: { buffer: output } }
      ]
    }));
    const encoder = device.createCommandEncoder();
    const pass = encoder.beginComputePass();
    pass.setPipeline(pipeline);
    pass.setBindGroup(0, group);
    pass.dispatchWorkgroups(1);
    pass.end();
    encoder.copyBufferToBuffer(output, 0, readback, 0, 24);
    await checked(device, '実行', () => device.queue.submit([encoder.finish()]));
    await readback.mapAsync(GPUMapMode.READ);
    const values = Array.from(new Float32Array(readback.getMappedRange()));
    readback.unmap();
    if (!lost) result.textContent = 'GPUの結果: ' + values.join(', ');
  } catch (error) {
    result.textContent = '実行できません: ' + error.message;
  } finally {
    finished = true;
    for (const value of buffers) value.destroy();
    device?.destroy();
  }
}
button.addEventListener('click', async () => {
  button.disabled = true;
  result.textContent = '処理中です。';
  try { await run(mode.value); }
  finally { button.disabled = false; }
});
</script></body></html>

処理は「CPUで配置する→uniformへ書く→シェーダーで六成分を読む→読取り用バッファへコピーする→JavaScriptで表示する」です。自動生成したlayoutを使い、この記事では接続番号の設計を固定しています。番号や用途で接続が止まる場合は、bind groupの照合で調べます。bind groupの作成時にはcreateBindGroupの条件に従い、必要な領域があるかも検証されます。

mapAsyncは、GPUの結果をJavaScriptが読める状態になるまで待つ処理です。mapAsyncの資料に従って読み取り後はunmapし、必要な値を普通の配列へコピーしてから表示しています。結果を読む前にバッファを破棄しない順序です。

領域が足りても、値の場所は間違えられる

「隙間なく詰める」は、32バイトの領域を確保したまま、二つ目の色を添字3へ書きます。最初の三成分は合いますが、GPUは二つ目を添字4から読むので、後半は0.25、0.5、0になります。0.75は一つ目の後ろの余白へ入り、読まれません。

この条件では接続の検証が通ります。バッファの大きさも用途も成立しており、間違っているのはアプリが決めた値の位置だからです。「GPUエラーがない」を「同じデータが届いた」へ読み替えないでください。

修正は、二つ目のdata.setの開始位置を4へ戻すことです。「16バイトごとに配置」を選び直し、後半が0.75、0.25、0.5へ戻るか確認します。先頭の三つだけを見ると、誤った配置のままでも合格してしまいます。

六成分分だけ確保すると、接続時に止まる

もう一つの失敗は、Float32Arrayを6要素、24バイトに縮めることです。シェーダーは二要素配列として32バイトの領域を要求するため、「領域を24バイトにする」では接続の検証エラーになります。この場合は値の読取りへ進みません。

修正では領域を8要素へ戻すことと、二つ目を添字4へ書くことを両方行います。大きさだけ直すと、前節の値のずれが残ります。必要なサイズと書込み位置を別々に照合する理由がここにあります。

WebGPUの検証エラーは、通常のtryとcatchだけで取得できるとは限りません。完成例はpushErrorScopeの使い方に沿ってvalidationを囲み、popErrorScopeの結果から失敗を知らせます。英文全体は実装により異なるので、「接続の段階で失敗した」という位置とサイズに関する内容を読みます。

この16バイトを、すべての型へ当てはめない

今回のstrideはvec3<f32>の配列から求めたものです。構造体の一つのvec3の後ろへf32を置く場合、そのf32が12バイト位置へ入ることもあります。配列の余白と構造体メンバーの配置を同じ規則で書き写さず、型全体の並びから計算します。行列や入れ子構造体の扱いは、この実験の範囲外です。

また、バッファの一部を接続するときのoffsetの制約は、内部の要素間隔とは別です。今回の接続offsetは0で固定しています。動的offset用の整列を調べて、その値で色ごとの間隔まで広げる必要はありません。

作成した三つのバッファと、この実験専用のdeviceは、一回の実行後に破棄します。GPU資源の寿命の記事と同様、途中失敗でも片付ける場所を残していますが、WebGLとWebGPUのAPIは別です。共有deviceを持つアプリへ組み込むなら、部品からdevice全体を破棄する設計にはしません。

最後に、正常なら六成分が入力と一致する、隙間なく詰めると後半だけが変わる、24バイトなら接続エラーになる、の三条件を確かめます。失敗の後に正常をもう一度実行し、同じ六成分へ戻れば比較は完了です。今回の検証はmacOS上のChromiumとSwiftShaderによるソフトウェア実行で、実機GPUの性能や描画品質の評価ではありません。canvasを扱う次の段階では、カリキュラムの描画分野とあわせて表示サイズや代替表示も確認してください。

TRY A LESSON

記事の内容を、課題で確かめる

前提になる表示サイズの練習 →

WebGL・3Dの発展課題 →

記事についてのお問い合わせ

記載内容へのご質問や訂正のご連絡は、hello@junkbranding.comまでお願いします。

← 記事一覧に戻る