From 7d5c90257996055852681685d8713c559c1ec812 Mon Sep 17 00:00:00 2001 From: Zack Buhman Date: Sun, 9 Aug 2026 14:22:34 -0500 Subject: [PATCH] compute render pipeline --- compute.js | 145 ++++++++++++++++++++++++++++++++++++++++++++++ compute.wgsl | 63 ++++---------------- font.js | 1 - game-of-life.wgsl | 56 ++++++++++++++++++ index.js | 26 ++++++--- 5 files changed, 231 insertions(+), 60 deletions(-) create mode 100644 compute.js create mode 100644 game-of-life.wgsl diff --git a/compute.js b/compute.js new file mode 100644 index 0000000..b4f68ef --- /dev/null +++ b/compute.js @@ -0,0 +1,145 @@ +import { getPath } from "./common.js"; + +class Compute { + constructor(device, canvasFormat, wgsl) + { + const label = "compute"; + this.label = label; + + ////////////////////////////////////////////////////////////////////// + // shader module + ////////////////////////////////////////////////////////////////////// + + const shaderModule = device.createShaderModule({ + label: `${label} shader`, + code: wgsl, + }); + + ////////////////////////////////////////////////////////////////////// + // bind group layout + ////////////////////////////////////////////////////////////////////// + + this.bindGroupLayout = { + compute: device.createBindGroupLayout({ + label: `${label} bind group layout`, + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + storageTexture: { + access: "write-only", + format: canvasFormat, + viewDimension: "2d", + } + } + ] + }), + }; + + ////////////////////////////////////////////////////////////////////// + // pipeline + ////////////////////////////////////////////////////////////////////// + + const pipelineLayout = device.createPipelineLayout({ + label: `${label} pipeline layout`, + bindGroupLayouts: [ + this.bindGroupLayout.compute, + ], + }); + + this.computePipeline = device.createComputePipeline({ + label: `${label} compute pipeline`, + layout: pipelineLayout, + compute: { + module: shaderModule, + }, + }); + + this.textures = undefined; + } + + recreateTextures(device, canvasTexture) + { + const label = this.label; + + if (this.textures !== undefined && this.textures[0].width == canvasTexture.width && this.textures[0].height == canvasTexture.height) { + return; + } + if (this.textures !== undefined) { + for (let i = 0; i < 2; i++) { + this.textures[i].destroy(); + } + } + + ////////////////////////////////////////////////////////////////////// + // textures + ////////////////////////////////////////////////////////////////////// + + this.textures = []; + for (let i = 0; i < 2; i++) { + const texture = device.createTexture({ + label: `${label} texture ${i}`, + format: canvasTexture.format, + usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_SRC, + size: [canvasTexture.width, canvasTexture.height], + }); + this.textures.push(texture); + } + + ////////////////////////////////////////////////////////////////////// + // bind group + ////////////////////////////////////////////////////////////////////// + + this.bindGroups = [] + for (let i = 0; i < 2; i++) { + const bindGroup = device.createBindGroup({ + label: `${label} bind group ${i}`, + layout: this.bindGroupLayout.compute, + entries: [{ + binding: 0, + resource: this.textures[i].createView(), + }], + }); + this.bindGroups.push(bindGroup); + } + } + + render(device, frameNumber, canvasTexture) + { + this.recreateTextures(device, canvasTexture); + + const label = this.label; + const encoder = device.createCommandEncoder({ label: `${label} command encoder` }); + const computePass = encoder.beginComputePass({ label: `${label} compute pass` }); + + computePass.setPipeline(this.computePipeline); + computePass.setBindGroup(0, this.bindGroups[frameNumber]); + computePass.dispatchWorkgroups(canvasTexture.width / 16, canvasTexture.height / 16); + computePass.end(); + + encoder.copyTextureToTexture({ + texture: this.textures[frameNumber], + }, { + texture: canvasTexture + }, { + width: canvasTexture.width, + height: canvasTexture.height, + depthOrArrayLayers: 1, + }); + + const commandBuffer = encoder.finish(); + device.queue.submit([commandBuffer]); + } +} + +async function loadCompute(device, canvasFormat) +{ + const canvasStorage = `@group(0) @binding(0) var out: texture_storage_2d<${canvasFormat}, write>;\n` + const wgsl = canvasStorage + await getPath("compute.wgsl"); + + const compute = new Compute(device, canvasFormat, wgsl); + + return compute; +} + +export { loadCompute }; diff --git a/compute.wgsl b/compute.wgsl index 851fbd8..379c263 100644 --- a/compute.wgsl +++ b/compute.wgsl @@ -1,56 +1,17 @@ -struct Configuration { - gridSize: vec2f, - aspect: vec2f, - matrix: mat4x4f -}; - -@group(0) @binding(0) var config: Configuration; -@group(0) @binding(1) var stateIn: array; -@group(0) @binding(2) var stateOut: array; - -fn cellIndex(cell: vec2i) -> u32 -{ - let x = (cell.x % i32(config.gridSize.x)); - let y = (cell.y % i32(config.gridSize.y)); - return u32(y * i32(config.gridSize.x) + x); -} - -fn cellActive(cell: vec2i) -> u32 -{ - return stateIn[cellIndex(cell)]; -} - -const directions: array = array(vec2i(-1, -1), - vec2i(-1, 0), - vec2i(-1, 1), - vec2i(0, -1), - //0, 0) - vec2i(0, 1), - vec2i(1, -1), - vec2i(1, 0), - vec2i(1, 1)); @compute -@workgroup_size(16) -fn computeMain(@builtin(global_invocation_id) globalInvocationId: vec3u) +@workgroup_size(16, 16) +fn computeMain(@builtin(global_invocation_id) id: vec3u) { - let cell: vec2i = vec2i(i32(globalInvocationId.x) % i32(config.gridSize.x), - i32(globalInvocationId.x) / i32(config.gridSize.y)); - - var activeNeighbors: u32 = 0; - for (var i: u32 = 0; i < 8; i++) { - activeNeighbors += cellActive(directions[i] + cell); + let outSize = textureDimensions(out); + if (id.x >= outSize.x || id.y >= outSize.y) { + return; } - let index = cellIndex(cell); - switch activeNeighbors { - case 2: { - stateOut[index] = stateIn[index]; - } - case 3: { - stateOut[index] = 1; - } - default: { - stateOut[index] = 0; - } - } + let coordinate = vec2f(f32(id.x) + 0.5, f32(outSize.y - id.y) - 0.5); + let outSizef = vec2f(outSize); + let uv = (coordinate * 2.0 - outSizef) / outSizef.y; + + let color = vec3f(uv.xy, 0); + + textureStore(out, id.xy, vec4f(color, 1.0)); } diff --git a/font.js b/font.js index f6cfb1b..6dcc92f 100644 --- a/font.js +++ b/font.js @@ -228,7 +228,6 @@ class FontRenderer { ], }); - console.log(canvasFormat); this.renderPipelines = {} for (let i = 1; i <= 4; i += 3) { this.renderPipelines[i] = device.createRenderPipeline({ diff --git a/game-of-life.wgsl b/game-of-life.wgsl new file mode 100644 index 0000000..851fbd8 --- /dev/null +++ b/game-of-life.wgsl @@ -0,0 +1,56 @@ +struct Configuration { + gridSize: vec2f, + aspect: vec2f, + matrix: mat4x4f +}; + +@group(0) @binding(0) var config: Configuration; +@group(0) @binding(1) var stateIn: array; +@group(0) @binding(2) var stateOut: array; + +fn cellIndex(cell: vec2i) -> u32 +{ + let x = (cell.x % i32(config.gridSize.x)); + let y = (cell.y % i32(config.gridSize.y)); + return u32(y * i32(config.gridSize.x) + x); +} + +fn cellActive(cell: vec2i) -> u32 +{ + return stateIn[cellIndex(cell)]; +} + +const directions: array = array(vec2i(-1, -1), + vec2i(-1, 0), + vec2i(-1, 1), + vec2i(0, -1), + //0, 0) + vec2i(0, 1), + vec2i(1, -1), + vec2i(1, 0), + vec2i(1, 1)); +@compute +@workgroup_size(16) +fn computeMain(@builtin(global_invocation_id) globalInvocationId: vec3u) +{ + let cell: vec2i = vec2i(i32(globalInvocationId.x) % i32(config.gridSize.x), + i32(globalInvocationId.x) / i32(config.gridSize.y)); + + var activeNeighbors: u32 = 0; + for (var i: u32 = 0; i < 8; i++) { + activeNeighbors += cellActive(directions[i] + cell); + } + + let index = cellIndex(cell); + switch activeNeighbors { + case 2: { + stateOut[index] = stateIn[index]; + } + case 3: { + stateOut[index] = 1; + } + default: { + stateOut[index] = 0; + } + } +} diff --git a/index.js b/index.js index 53ed6cd..319e7b5 100644 --- a/index.js +++ b/index.js @@ -4,6 +4,7 @@ import { loadLight } from "./light.js"; import { loadFlat } from "./flat.js"; import { loadSnake } from "./snake.js"; import { loadFont } from "./font.js"; +import { loadCompute } from "./compute.js"; if (!navigator.gpu) { throw new Error("WebGPU not supported on this browser."); @@ -14,19 +15,24 @@ if (!adapter) { throw new Error("No WebGPU adapter"); } -const device = await adapter.requestDevice(); +const hasBgraStorage = adapter.features.has("bgra8unorm-storage"); +const device = await adapter.requestDevice({ + requiredFeatures: hasBgraStorage ? ["bgra8unorm-storage"] : [] +}); const canvas = document.querySelector("canvas"); canvas.width = canvas.clientWidth canvas.height = canvas.clientHeight const context = canvas.getContext("webgpu"); -const canvasFormat = navigator.gpu.getPreferredCanvasFormat(); +const canvasFormat = hasBgraStorage ? navigator.gpu.getPreferredCanvasFormat() : "rgba8unorm"; context.configure({ device: device, format: canvasFormat, alphaMode: 'premultiplied', + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT }); +console.log("canvasFormat", canvasFormat); const encoder = device.createCommandEncoder(); @@ -151,7 +157,8 @@ module.instance.exports.camera_init(cameraStateAddress); //const lightRenderer = await loadLight(device, canvasFormat, viewUniformBuffer); //const flatRenderer = await loadFlat(device, canvasFormat, viewUniformBuffer, module); //const snakeRenderer = await loadSnake(device, canvasFormat, viewUniformBuffer, module); -const fontRenderer = await loadFont(device, canvasFormat, viewUniformBuffer, module); +//const fontRenderer = await loadFont(device, canvasFormat, viewUniformBuffer, module); +const compute = await loadCompute(device, canvasFormat, module); const KEY = { A: 65, @@ -333,8 +340,9 @@ function render2() canvas.width = canvas.clientWidth * window.devicePixelRatio; canvas.height = canvas.clientHeight * window.devicePixelRatio; - const configuration = fontRenderer.update(device, frameNumber); - const sampleCount = configuration.multisampleCount; + //const configuration = fontRenderer.update(device, frameNumber); + //const sampleCount = configuration.multisampleCount; + const sampleCount = 1; //const configuration = updateSliders(flatSliders); //const configuration = updateSliders(snakeSliders); @@ -380,7 +388,7 @@ function render2() //snakeRenderer.update(device, frameNumber, mousePosition, configuration); //snakeRenderer.render(renderPass, frameNumber); - fontRenderer.render(renderPass, frameNumber, sampleCount); + //fontRenderer.render(renderPass, frameNumber, sampleCount); renderPass.end(); @@ -397,8 +405,10 @@ function render2() }); } - const commandBuffer = encoder.finish(); - device.queue.submit([commandBuffer]); + //const commandBuffer = encoder.finish(); + //device.queue.submit([commandBuffer]); + + compute.render(device, frameNumber, canvasTexture); frameNumber = (frameNumber + 1) % 2;