この記事の目次
シェーダー側ではbindingを2にしたのに、JavaScript側では配列の先頭だから0を書いてしまう。WebGPUの資源接続では、配列の位置と接続番号を区別する必要があります。番号が合っていても、バッファを作ったときの用途が違えば接続できません。
この練習は、JavaScriptの非同期関数とオブジェクト、WGSLの宣言を読める人向けです。入力0.5をGPUで二倍にし、結果1を文章へ戻します。WebGPUのcompute shaderを使うので、canvas、頂点データ、動きの制御はありません。描画の前に、資源へ届く経路を調べます。
groupとbindingを、二段階の住所として読む
bind groupは、複数のバッファなどをまとめてシェーダーへ渡す仕組みです。layoutは受け取る資源の番号と用途を定め、bind groupはそこへ実際のバッファを対応させます。createBindGroupLayoutの資料では、番号、利用するシェーダー段階、バッファ種別を別々に指定します。
今回の住所は、入力がgroup 0のbinding 2、出力がgroup 0のbinding 5です。0から連続するbinding番号にそろえる必要はありません。
| 対象 | WGSLの宣言 | JavaScript側の対応 |
|---|---|---|
| グループ全体 | @group(0) | bindGroupLayoutsの0番目、setBindGroupの第1引数0 |
| 読取り専用の入力 | @binding(2)、var<uniform> |
binding: 2、buffer.type: uniform |
| 計算結果の出力 | @binding(5)、var<storage, read_write> |
binding: 5、buffer.type: storage |
| 計算する段階 | @compute | visibility: GPUShaderStage.COMPUTE |
WebGPU仕様の資源接続にもあるように、groupの番号とグループ内のbinding番号は別です。setBindGroupへ2を渡しても、binding 2を選ぶ意味にはなりません。今回はグループを一つだけ使います。
先頭に出力を書いても、番号が合えば接続できる
完成例のbind groupのentriesは、binding 5を先、binding 2を後に書いてあります。それでも入力と出力を取り違えません。createBindGroupの説明のとおり、照合するのは各entryのbindingです。
次をbinding.htmlへ保存します。Python 3があるフォルダでpython3 -m http.server 8768 --bind 127.0.0.1を実行し、http://127.0.0.1:8768/binding.htmlを開いてください。追加ライブラリは不要で、終了はCtrl+Cです。WebGPUを利用できるブラウザーと、HTTPSまたは信頼されるローカル環境を使います。
<!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>入力0.5を二倍にし、結果1を読み戻す実験です。</p>
<label for="mode">試す条件</label>
<select id="mode">
<option value="ok">正しい接続</option>
<option value="number">binding番号を変える</option>
<option value="usage">bufferの用途を変える</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 layout = device.createBindGroupLayout({ entries: [
{ binding: 2, visibility: GPUShaderStage.COMPUTE,
buffer: { type: 'uniform', minBindingSize: 16 } },
{ binding: 5, visibility: GPUShaderStage.COMPUTE,
buffer: { type: 'storage', minBindingSize: 4 } }
] });
const module = device.createShaderModule({ code: `
@group(0) @binding(2) var<uniform> settings: vec4<f32>;
@group(0) @binding(5) var<storage, read_write> result: f32;
@compute @workgroup_size(1) fn main() {
result = settings.x * 2.0;
}` });
const pipeline = await device.createComputePipelineAsync({
layout: device.createPipelineLayout({ bindGroupLayouts: [layout] }),
compute: { module, entryPoint: 'main' }
});
const inputUsage = kind === 'usage' ? GPUBufferUsage.STORAGE : GPUBufferUsage.UNIFORM;
const input = buffer(16, inputUsage | GPUBufferUsage.COPY_DST);
const output = buffer(4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
const readback = buffer(4, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
device.queue.writeBuffer(input, 0, new Float32Array([0.5, 0, 0, 0]));
const group = await checked(device, '接続', () => device.createBindGroup({
layout,
entries: [
{ binding: 5, resource: { buffer: output } },
{ binding: kind === 'number' ? 3 : 2, resource: { buffer: input } }
]
}));
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, 4);
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>
入力のvec4<f32>は16バイト、結果のf32は4バイトです。minBindingSizeへそれぞれの必要量を明記し、pipelineのlayoutとシェーダーの契約を作成時にも照合できる構成にしています。これはバッファを確保する指定ではなく、接続を許す最小サイズの指定です。
writeBufferで0.5を書き、compute passで入力を二倍にします。出力は読取り専用のコピー先へ移し、mapAsyncの完了後にJavaScriptへ取り出します。「接続できた」だけで終えず、値が結果1になるところまで確認できます。
bindingを3へ変えたときは、番号の表へ戻る
「binding番号を変える」を選ぶと、入力entryだけがbinding 3になります。layoutが要求しているのは2と5なので、createBindGroupの検証で止まるのが期待結果です。配列の二番目へ置いたことや、buffer変数に正しい値があることでは補えません。
修正は、入力entryをbinding 2へ戻すことです。WGSL、layout、bind groupの三か所を表で縦に照合します。今回は前二つを固定しているため最後だけを直しますが、実際の変更で三か所ともずれていれば、仕様として選んだ一つの番号へそろえます。
「正しい接続」へ戻して結果1を確認してください。entriesを逆順に書き換えても1なら、配列の位置を接続番号と取り違えていないことを追加で確かめられます。正常時に順番を変える実験と、番号を変える失敗を混ぜずに試します。
番号が正しくても、usageは後から付け替えられない
「bufferの用途を変える」では番号を2へ戻したまま、入力バッファをSTORAGEとCOPY_DSTで作成します。layoutはuniformを要求するため、この条件も接続時に失敗します。値の0.5やバッファの16バイトを直しても、用途の不一致は解消しません。
修正はcreateBufferのusageへUNIFORMを指定してバッファを作り直すことです。既存のバッファへ別のusageを後付けする操作ではありません。COPY_DSTはwriteBufferで書き込むために併記しています。シェーダーの読取り用途と、CPUから更新する用途を両方満たします。
反対に、計算結果を書く出力をuniformへ変更してはいけません。WGSLのvar<storage, read_write>と、layoutのstorage、作成時のSTORAGEが対応します。何となくすべてのusageを足して通すより、入力と出力で何をするかを決め、その用途だけを付けます。
検証エラーはerror scopeで取り出します。通常のJavaScript例外と分けて取得した後、この記事のchecked関数が例外へ変えて表示しています。catchが直接すべてのGPUエラーを捕まえている、という読み方はしないでください。
自動layoutを使う場合と、明示する場合
一つのpipelineへ資源を渡す小さな処理なら、layout: autoとgetBindGroupLayoutで必要なlayoutを取得する方法もあります。今回明示したのは、受け渡す契約を表とコードで比較するためです。複数のpipelineで同じ資源配置を使う場合も、明示的なlayoutが候補になります。
自動layoutから作ったbind groupを、別pipelineでもそのまま使えると考えないでください。また明示したlayoutでも、シェーダー段階や型が変われば照合し直します。古い表を残したままシェーダーだけ更新すると、どの接続を意図したのか分からなくなります。
WebGLのシェーダー診断がプログラム同士の接続を調べるのに対し、ここで照合したのはJavaScriptが持つ実資源とWGSLの受取口です。接続に成功しても数値がずれる場合は、uniform内部の配置を別に調べます。WGSLの構文エラーならシェーダーやpipeline作成の段階へ戻り、bind groupの番号だけを変え続けないことも切り分けになります。
結果1と、二種類の接続失敗を記録する
確認条件は、正常時に結果1、番号変更で接続エラー、用途変更でも接続エラー、そして正常へ戻すと再び1、です。エラーの英文全体は実装差があるため、関係するbinding番号やバッファ用途を読みます。キーボードで条件を選び、ボタンを押せることも確認します。
APIがない、adapterがない、deviceを取得できない場合は計算をせず文章へ戻ります。作成したバッファと実験専用deviceはfinallyで破棄します。デバイス喪失後の自動復元は含めず、再実行で取得からやり直します。共有のdeviceを使うアプリでは、この後始末を部品の所有範囲へ変更してください。
今回の動作確認はmacOS上のChromiumとSwiftShaderによるソフトウェアGPUです。実機の速度や他ブラウザーの対応を測った結果ではありません。文章出力なのでDPR・canvasのリサイズは対象外です。JavaScriptの無料レッスンから非同期処理へ進む場合も、「接続を作れた」と「期待する数値を得た」を二つの確認結果として残せます。
TRY A LESSON
記事の内容を、課題で確かめる
記載内容へのご質問や訂正のご連絡は、hello@junkbranding.comまでお願いします。