maku693です。この記事は はてなエンジニア Advent Calendar 2025の11日目の記事です。
以前 WebGPUの紹介 - The Third Law という記事を書いてから6年が経ち、WebGPUが最新のブラウザで広くサポートされるようになりました*1。この記事では、改めて最新のWebGPUについてまとめてみます。
WebGPUとは
WebGPU*2とは、W3Cによって策定されている、Web上でGPUを活用するためのAPIです。比較的最近のGPUのアーキテクチャに合わせたAPIを持ち、WebGLに比べてパフォーマンスがよかったり、GPGPUが最初からサポートされていたりします。DirectX 12やVulkan, MetalといったAPIに似ていますが、ブラウザ上に実装されることを前提としているため、セキュリティやプライバシーに配慮した仕様となっていたり、他のWeb APIとの連携が考慮されているのが特徴です。
前掲の記事を書いた時点の仕様はWorking Draftで、ブラウザごとに細かなAPIやシェーディング言語がバラバラでしたが、現在はCandidate Recommendation Draft (CRD) となっており、ほぼ安定しているとみなしていいでしょう。合わせて利用されるWebGPU Shading Language*3についても同様にCRDで、どちらも現在主要ブラウザの最新版で広く利用できます。
WebGPUに触れてみる 2025
今回もレンダリングとGPGPUのデモとして、以前古いAPIで書いたパーティクルの実装を移植(+ちょっと気になった部分を書き換え)してみました。動作するデモは以下です。色々なブラウザで動いて嬉しいですね。
リポジトリはこちら。
移植にあたって必要だった差分は以下です。
Diff
diff --git a/script.js b/script.js index c8a7e95..16b0f92 100644 --- a/script.js +++ b/script.js @@ -2,24 +2,24 @@ const particleCount = 4096; const shaderSource = ` struct Particle { - float2 position; - float2 velocity; + position: vec2f, + velocity: vec2f, } -[numthreads(1, 1, 1)] -compute void ComputeMain( - device Particle[] prevParticles : register(u0), - device Particle[] nextParticles : register(u1), - float3 threadID : SV_DispatchThreadID -) { - uint index = uint(threadID.x); +@group(0) @binding(0) var<storage, read_write> prevParticles: array<Particle>; +@group(0) @binding(1) var<storage, read_write> nextParticles: array<Particle>; +@compute @workgroup_size(1) +fn ComputeMain( + @builtin(global_invocation_id) threadID: vec3u, +) { + var index = threadID.x; if (${particleCount} <= index) { return; } - float2 position = prevParticles[index].position; - float2 velocity = prevParticles[index].velocity; + var position = prevParticles[index].position; + var velocity = prevParticles[index].velocity; position += velocity; @@ -41,58 +41,57 @@ compute void ComputeMain( } struct VertexOut { - float4 position : SV_Position; - float pointSize : PSIZE; + @builtin(position) position: vec4f, } -vertex VertexOut VertexMain(float2 position : attribute(0)) { - VertexOut out; - out.position = float4(position, 0, 1); - out.pointSize = 2; +@vertex +fn VertexMain(@location(0) position: vec2f) -> VertexOut { + var out: VertexOut; + out.position = vec4f(position, 0, 1); return out; } -fragment float4 FragmentMain() : SV_Target 0 { - return float4(1, 1, 1, 1); +@fragment +fn FragmentMain() -> @location(0) vec4f { + return vec4f(1, 1, 1, 1); } `; async function main() { if (!navigator.gpu) { document.body.innerText = - "Sorry, your browser isn't supported. Please use Safari TP."; + "Sorry, your browser isn't supported. Please use latest browser version."; return; } // general resources const canvas = document.getElementById("canvas"); - const context = canvas.getContext("gpu"); + const context = canvas.getContext("webgpu"); const adapter = await navigator.gpu.requestAdapter(); const device = await adapter.requestDevice(); - const swapChain = context.configureSwapChain({ + context.configure({ device, - format: "bgra8unorm", + format: navigator.gpu.getPreferredCanvasFormat(), }); const shaderModule = device.createShaderModule({ code: shaderSource, - isWHLSL: true, }); // compute pipeline const computeBindGroupLayout = device.createBindGroupLayout({ - bindings: [ + entries: [ { binding: 0, visibility: GPUShaderStage.COMPUTE, - type: "storage-buffer", + buffer: { type: "storage" }, }, { binding: 1, visibility: GPUShaderStage.COMPUTE, - type: "storage-buffer", + buffer: { type: "storage" }, }, ], }); @@ -101,7 +100,7 @@ async function main() { }); const computePipeline = device.createComputePipeline({ layout: computePipelineLayout, - computeStage: { + compute: { module: shaderModule, entryPoint: "ComputeMain", }, @@ -109,31 +108,16 @@ async function main() { // render pipeline const renderPipeline = device.createRenderPipeline({ - vertexStage: { + vertex: { module: shaderModule, entryPoint: "VertexMain", - }, - fragmentStage: { - module: shaderModule, - entryPoint: "FragmentMain", - }, - primitiveTopology: "point-list", - colorStates: [ - { - format: "bgra8unorm", - alphaBlend: {}, - colorBlend: {}, - }, - ], - vertexInput: { - indexFormat: "uint32", - vertexBuffers: [ + buffers: [ { - stride: 8 * 2, + arrayStride: 8 * 2, stepMode: "vertex", - attributeSet: [ + attributes: [ { - format: "float2", + format: "float32x2", offset: 0, shaderLocation: 0, }, @@ -141,6 +125,17 @@ async function main() { }, ], }, + fragment: { + module: shaderModule, + entryPoint: "FragmentMain", + targets: [ + { + format: navigator.gpu.getPreferredCanvasFormat(), + }, + ], + }, + primitive: { topology: "point-list" }, + layout: "auto", }); // particle data @@ -154,11 +149,12 @@ async function main() { const particleBuffers = new Array(2); for (let i = 0; i < 2; i++) { - const [particleBuffer, particleArrayBuffer] = - await device.createBufferMapped({ - size: particleData.byteLength, - usage: GPUBufferUsage.VERTEX | GPUBufferUsage.STORAGE, - }); + const particleBuffer = await device.createBuffer({ + size: particleData.byteLength, + usage: GPUBufferUsage.VERTEX | GPUBufferUsage.STORAGE, + mappedAtCreation: true, + }); + const particleArrayBuffer = particleBuffer.getMappedRange(); new Float32Array(particleArrayBuffer).set(particleData); particleBuffer.unmap(); particleBuffers[i] = particleBuffer; @@ -168,22 +164,14 @@ async function main() { for (let i = 0; i < 2; i++) { const particleBindGroup = device.createBindGroup({ layout: computeBindGroupLayout, - bindings: [ + entries: [ { binding: 0, - resource: { - buffer: particleBuffers[i], - offset: 0, - size: particleData.byteLength, - }, + resource: particleBuffers[i], }, { binding: 1, - resource: { - buffer: particleBuffers[(i + 1) % 2], - offset: 0, - size: particleData.byteLength, - }, + resource: particleBuffers[(i + 1) % 2], }, ], }); @@ -197,25 +185,25 @@ async function main() { const computePassEncoder = commandEncoder.beginComputePass(); computePassEncoder.setPipeline(computePipeline); computePassEncoder.setBindGroup(0, particleBindGroups[t % 2]); - computePassEncoder.dispatch(particleCount, 1, 1); - computePassEncoder.endPass(); + computePassEncoder.dispatchWorkgroups(particleCount, 1, 1); + computePassEncoder.end(); const renderPassEncoder = commandEncoder.beginRenderPass({ colorAttachments: [ { - attachment: swapChain.getCurrentTexture().createDefaultView(), + view: context.getCurrentTexture(), loadOp: "clear", - clearColor: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 }, + clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 }, storeOp: "store", }, ], }); renderPassEncoder.setPipeline(renderPipeline); - renderPassEncoder.setVertexBuffers(0, [particleBuffers[(t + 1) % 2]], [0]); + renderPassEncoder.setVertexBuffer(0, particleBuffers[(t + 1) % 2]); renderPassEncoder.draw(particleCount, 1, 0, 0); - renderPassEncoder.endPass(); + renderPassEncoder.end(); - device.getQueue().submit([commandEncoder.finish()]); + device.queue.submit([commandEncoder.finish()]); t++; requestAnimationFrame(render);
改めて触ってみた印象としては、6年前と比べると、シェーダの文法が完全に置き換わっているのが目立ちます。これはWebGPU Shading Language (WGSL) という言語です。文法はOpenGL系のGLSL・DirectX系のHLSLのどちらとも違っており、どことなくRustやTypeScriptっぽさを感じる新しいものです。今回初めて書いたのですが、地味に引数や構造体のフィールド定義のtrailing commaが許容されているのが気に入りました。
新しいシェーディング言語が採用されることになった経緯は仕様には特に書かれていないのですが、ざっくり調べた限りでは、既存の言語やIRだとセキュリティや知的財産権の懸念があり、新たにテキスト形式の言語を用意することになったようです。WebGPUに合わせて新たに作られただけあって、APIとのインタフェースの齟齬が少なく、書く側としては嬉しいですね。個人的にはGLSLより読み書きしやすいと感じました。
また、WebGPU APIそのものについては、プロパティ名やメソッド名、メソッドの引数の構造が整理されて、全体的に冗長な記述をしなくて済むようになっています。初期のドラフト実装の時点でもわかりやすいAPIでしたが、それにさらに磨きがかかっている印象です。一方、pointを描画する際のサイズの指定ができなくなっていることにも気づきました。あまり使う機会はないと思いますが、1px以上の大きさの点を描画したい場合は自前で実装する必要がありそうです。じっくりとは確認していませんが、これ以外にもプラットフォーム固有の機能などは使えなくなっているかもしれません。また、Chromeでは動くがSafariでは動かないショートハンドがある*4など、マイナーな互換性の問題はまだ多少ありそうです。
今回はAPIの差分を見たかったので直接WebGPUを呼び出しましたが、エコシステムも以前に比べて充実してきていそうです。Three.jsやBabylon.jsなどの著名ライブラリで既にサポートされているのはもちろん、TypeGPU*5というTypeScriptを使ってWebGPU側のbufferの定義とJS側のArrayBufferの関連付けをサポートしてくれるライブラリなども登場してきています(この作業は慎重さが必要で正直面倒なのでライブラリが解決してくれるのは非常にありがたい)。ライブラリだけでなく、日本語の初心者向けの情報も増えてきており、WebGPUを使ったり入門したりしやすい環境が整ってきていると思います。
以上、現状のWebGPUに簡単に触れてみました。これを機にGPUやシェーディング言語に興味を持ってもらえれば幸いです。