From 9ad655ca2efc79f63f8e9ca220902e3bcc9e5d4a Mon Sep 17 00:00:00 2001 From: Rezmason Date: Thu, 11 Nov 2021 09:18:32 -0800 Subject: [PATCH] Created a pass-through post processing compute pass. The other post-processing passes will be changed over to this kind of thing. makePassFBO has now been split into makeRenderTarget and makeComputeTarget. --- js/webgpu/bloomPass.js | 32 +++++++------- js/webgpu/endPass.js | 2 +- js/webgpu/imagePass.js | 4 +- js/webgpu/main.js | 3 +- js/webgpu/palettePass.js | 4 +- js/webgpu/postProcessingPass.js | 52 ++++++++++++++++++++++ js/webgpu/rainPass.js | 8 ++-- js/webgpu/resurrectionPass.js | 4 +- js/webgpu/stripePass.js | 4 +- js/webgpu/utils.js | 11 ++++- shaders/wgsl/postProcessingPass.wgsl | 66 ++++++++++++++++++++++++++++ 11 files changed, 158 insertions(+), 32 deletions(-) create mode 100644 js/webgpu/postProcessingPass.js create mode 100644 shaders/wgsl/postProcessingPass.wgsl diff --git a/js/webgpu/bloomPass.js b/js/webgpu/bloomPass.js index d96dd34..fcda98e 100644 --- a/js/webgpu/bloomPass.js +++ b/js/webgpu/bloomPass.js @@ -1,5 +1,5 @@ import { structs } from "/lib/gpu-buffer.js"; -import { loadShader, makeUniformBuffer, makeBindGroup, makePassFBO, makePass } from "./utils.js"; +import { loadShader, makeUniformBuffer, makeBindGroup, makeRenderTarget, makePass } from "./utils.js"; // The bloom pass is basically an added blur of the high-pass rendered output. // The blur approximation is the sum of a pyramid of downscaled textures. @@ -17,7 +17,7 @@ export default (context, getInputs) => { // If there's no bloom to apply, return a no-op pass with an empty bloom texture if (!enabled) { - const emptyTexture = makePassFBO(device, 1, 1, canvasFormat); + const emptyTexture = makeRenderTarget(device, 1, 1, canvasFormat); const getOutputs = () => ({ ...getInputs(), bloom: emptyTexture }); return makePass(getOutputs); } @@ -26,8 +26,8 @@ export default (context, getInputs) => { // TODO: generate sum shader code - const fbo = makePassFBO(device, 1, 1, canvasFormat); - const getOutputs = () => ({ ...getInputs(), bloom: fbo }); // TODO + const renderTarget = makeRenderTarget(device, 1, 1, canvasFormat); + const getOutputs = () => ({ ...getInputs(), bloom: renderTarget }); // TODO let blurRenderPipeline; let sumRenderPipeline; @@ -69,23 +69,23 @@ export default (context, getInputs) => { /* -// A pyramid is just an array of FBOs, where each FBO is half the width -// and half the height of the FBO below it. +// A pyramid is just an array of Targets, where each Target is half the width +// and half the height of the Target below it. const makePyramid = (regl, height, halfFloat) => Array(height) .fill() - .map((_) => makePassFBO(regl, halfFloat)); + .map((_) => makeRenderTarget(regl, halfFloat)); const resizePyramid = (pyramid, vw, vh, scale) => pyramid.forEach((fbo, index) => fbo.resize(Math.floor((vw * scale) / 2 ** index), Math.floor((vh * scale) / 2 ** index))); export default ({ regl, config }, inputs) => { - // Build three pyramids of FBOs, one for each step in the process + // Build three pyramids of Targets, one for each step in the process const highPassPyramid = makePyramid(regl, pyramidHeight, config.useHalfFloat); const hBlurPyramid = makePyramid(regl, pyramidHeight, config.useHalfFloat); const vBlurPyramid = makePyramid(regl, pyramidHeight, config.useHalfFloat); - const output = makePassFBO(regl, config.useHalfFloat); + const output = makeRenderTarget(regl, config.useHalfFloat); // The high pass restricts the blur to bright things in our input texture. const highPassFrag = loadText("shaders/glsl/highPass.frag.glsl"); @@ -99,7 +99,7 @@ export default ({ regl, config }, inputs) => { }); // A 2D gaussian blur is just a 1D blur done horizontally, then done vertically. - // The FBO pyramid's levels represent separate levels of detail; + // The Target pyramid's levels represent separate levels of detail; // by blurring them all, this basic blur approximates a more complex gaussian: // https://web.archive.org/web/20191124072602/https://software.intel.com/en-us/articles/compute-shader-hdr-and-bloom @@ -150,12 +150,12 @@ export default ({ regl, config }, inputs) => { }, () => { for (let i = 0; i < pyramidHeight; i++) { - const highPassFBO = highPassPyramid[i]; - const hBlurFBO = hBlurPyramid[i]; - const vBlurFBO = vBlurPyramid[i]; - highPass({ fbo: highPassFBO, frag: highPassFrag.text(), tex: inputs.primary }); - blur({ fbo: hBlurFBO, frag: blurFrag.text(), tex: highPassFBO, direction: [1, 0] }); - blur({ fbo: vBlurFBO, frag: blurFrag.text(), tex: hBlurFBO, direction: [0, 1] }); + const highPassTarget = highPassPyramid[i]; + const hBlurTarget = hBlurPyramid[i]; + const vBlurTarget = vBlurPyramid[i]; + highPass({ fbo: highPassTarget, frag: highPassFrag.text(), tex: inputs.primary }); + blur({ fbo: hBlurTarget, frag: blurFrag.text(), tex: highPassTarget, direction: [1, 0] }); + blur({ fbo: vBlurTarget, frag: blurFrag.text(), tex: hBlurTarget, direction: [0, 1] }); } sumPyramid(); diff --git a/js/webgpu/endPass.js b/js/webgpu/endPass.js index 5d3f99e..ad62d8d 100644 --- a/js/webgpu/endPass.js +++ b/js/webgpu/endPass.js @@ -1,4 +1,4 @@ -import { loadShader, makeBindGroup, makePassFBO, makePass } from "./utils.js"; +import { loadShader, makeBindGroup, makePass } from "./utils.js"; const numVerticesPerQuad = 2 * 3; diff --git a/js/webgpu/imagePass.js b/js/webgpu/imagePass.js index 0369a01..e3ed30f 100644 --- a/js/webgpu/imagePass.js +++ b/js/webgpu/imagePass.js @@ -1,4 +1,4 @@ -import { loadTexture, loadShader, makeBindGroup, makePassFBO, makePass } from "./utils.js"; +import { loadTexture, loadShader, makeBindGroup, makeRenderTarget, makePass } from "./utils.js"; // Multiplies the rendered rain and bloom by a loaded in image @@ -58,7 +58,7 @@ export default (context, getInputs) => { const setSize = (width, height) => { output?.destroy(); - output = makePassFBO(device, width, height, canvasFormat); + output = makeRenderTarget(device, width, height, canvasFormat); }; const execute = (encoder) => { diff --git a/js/webgpu/main.js b/js/webgpu/main.js index e919698..021bdba 100644 --- a/js/webgpu/main.js +++ b/js/webgpu/main.js @@ -7,6 +7,7 @@ import makePalettePass from "./palettePass.js"; import makeStripePass from "./stripePass.js"; import makeImagePass from "./imagePass.js"; import makeResurrectionPass from "./resurrectionPass.js"; +import makePostProcessingPass from "./postProcessingPass.js"; import makeEndPass from "./endPass.js"; const effects = { @@ -52,7 +53,7 @@ export default async (canvas, config) => { }; const effectName = config.effect in effects ? config.effect : "plain"; - const pipeline = makePipeline(context, [makeRain, makeBloomPass, effects[effectName], makeEndPass]); + const pipeline = makePipeline(context, [makeRain, makeBloomPass, effects[effectName], makePostProcessingPass, makeEndPass]); await Promise.all(pipeline.map((step) => step.ready)); diff --git a/js/webgpu/palettePass.js b/js/webgpu/palettePass.js index 1c8918e..94f6f41 100644 --- a/js/webgpu/palettePass.js +++ b/js/webgpu/palettePass.js @@ -1,5 +1,5 @@ import { structs } from "/lib/gpu-buffer.js"; -import { loadShader, makeUniformBuffer, makeBindGroup, makePassFBO, makePass } from "./utils.js"; +import { loadShader, makeUniformBuffer, makeBindGroup, makeRenderTarget, makePass } from "./utils.js"; // Maps the brightness of the rendered rain and bloom to colors // in a linear gradient buffer generated from the passed-in color sequence @@ -135,7 +135,7 @@ export default (context, getInputs) => { const setSize = (width, height) => { output?.destroy(); - output = makePassFBO(device, width, height, canvasFormat); + output = makeRenderTarget(device, width, height, canvasFormat); }; const execute = (encoder) => { diff --git a/js/webgpu/postProcessingPass.js b/js/webgpu/postProcessingPass.js new file mode 100644 index 0000000..cb9deb8 --- /dev/null +++ b/js/webgpu/postProcessingPass.js @@ -0,0 +1,52 @@ +import { structs, byteSizeOf } from "/lib/gpu-buffer.js"; +import { makeComputeTarget, loadShader, makeUniformBuffer, makeBindGroup, makePass } from "./utils.js"; + +export default (context, getInputs) => { + const { config, device, timeBuffer } = context; + + const assets = [loadShader(device, "shaders/wgsl/postProcessingPass.wgsl")]; + + let configBuffer; + let computePipeline; + let output; + let screenSize; + + const getOutputs = () => ({ + primary: output, + }); + + const ready = (async () => { + const [postProcessingShader] = await Promise.all(assets); + + computePipeline = device.createComputePipeline({ + compute: { + module: postProcessingShader.module, + entryPoint: "computeMain", + }, + }); + + const configUniforms = structs.from(postProcessingShader.code).Config; + configBuffer = makeUniformBuffer(device, configUniforms, { + /* TODO */ + }); + })(); + + const setSize = (width, height) => { + output?.destroy(); + output = makeComputeTarget(device, width, height); + screenSize = [width, height]; + }; + + const execute = (encoder) => { + const inputs = getInputs(); + const tex = inputs.primary; + const computePass = encoder.beginComputePass(); + computePass.setPipeline(computePipeline); + const computeBindGroup = makeBindGroup(device, computePipeline, 0, [configBuffer, timeBuffer, tex.createView(), output.createView()]); + computePass.setBindGroup(0, computeBindGroup); + computePass.dispatch(Math.ceil(screenSize[0] / 32), screenSize[1], 1); + computePass.endPass(); + }; + + return makePass(getOutputs, ready, setSize, execute); +}; diff --git a/js/webgpu/rainPass.js b/js/webgpu/rainPass.js index 661503b..fe991b8 100644 --- a/js/webgpu/rainPass.js +++ b/js/webgpu/rainPass.js @@ -1,5 +1,5 @@ import { structs, byteSizeOf } from "/lib/gpu-buffer.js"; -import { makePassFBO, loadTexture, loadShader, makeUniformBuffer, makeBindGroup, makePass } from "./utils.js"; +import { makeRenderTarget, loadTexture, loadShader, makeUniformBuffer, makeBindGroup, makePass } from "./utils.js"; const { mat4, vec3 } = glMatrix; @@ -167,14 +167,14 @@ export default (context, getInputs) => { // Update output?.destroy(); - output = makePassFBO(device, width, height, canvasFormat); + output = makeRenderTarget(device, width, height, canvasFormat); highPassOutput?.destroy(); - highPassOutput = makePassFBO(device, width, height, canvasFormat); + highPassOutput = makeRenderTarget(device, width, height, canvasFormat); }; const execute = (encoder) => { - // We render the code into an FBO using MSDFs: https://github.com/Chlumsky/msdfgen + // We render the code into an Target using MSDFs: https://github.com/Chlumsky/msdfgen const computePass = encoder.beginComputePass(); computePass.setPipeline(computePipeline); diff --git a/js/webgpu/resurrectionPass.js b/js/webgpu/resurrectionPass.js index a54219f..91569b7 100644 --- a/js/webgpu/resurrectionPass.js +++ b/js/webgpu/resurrectionPass.js @@ -1,5 +1,5 @@ import { structs } from "/lib/gpu-buffer.js"; -import { loadShader, makeUniformBuffer, makePassFBO, makePass } from "./utils.js"; +import { loadShader, makeUniformBuffer, makeRenderTarget, makePass } from "./utils.js"; // Matrix Resurrections isn't in theaters yet, // and this version of the effect is still a WIP. @@ -60,7 +60,7 @@ export default (context, getInputs) => { const setSize = (width, height) => { output?.destroy(); - output = makePassFBO(device, width, height, canvasFormat); + output = makeRenderTarget(device, width, height, canvasFormat); }; const getOutputs = () => ({ diff --git a/js/webgpu/stripePass.js b/js/webgpu/stripePass.js index 4d29d41..ad23d5e 100644 --- a/js/webgpu/stripePass.js +++ b/js/webgpu/stripePass.js @@ -1,5 +1,5 @@ import { structs } from "/lib/gpu-buffer.js"; -import { loadShader, make1DTexture, makeUniformBuffer, makeBindGroup, makePassFBO, makePass } from "./utils.js"; +import { loadShader, make1DTexture, makeUniformBuffer, makeBindGroup, makeRenderTarget, makePass } from "./utils.js"; // Multiplies the rendered rain and bloom by a 1D gradient texture // generated from the passed-in color sequence @@ -96,7 +96,7 @@ export default (context, getInputs) => { const setSize = (width, height) => { output?.destroy(); - output = makePassFBO(device, width, height, canvasFormat); + output = makeRenderTarget(device, width, height, canvasFormat); }; const getOutputs = () => ({ diff --git a/js/webgpu/utils.js b/js/webgpu/utils.js index 9f69547..cb2dfff 100644 --- a/js/webgpu/utils.js +++ b/js/webgpu/utils.js @@ -27,13 +27,20 @@ const loadTexture = async (device, url) => { return texture; }; -const makePassFBO = (device, width, height, format = "bgra8unorm") => +const makeRenderTarget = (device, width, height, format) => device.createTexture({ size: [width, height, 1], format, usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST | GPUTextureUsage.RENDER_ATTACHMENT, }); +const makeComputeTarget = (device, width, height) => + device.createTexture({ + size: [width, height, 1], + format: "rgba8unorm", + usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, + }); + const loadShader = async (device, url) => { const response = await fetch(url); const code = await response.text(); @@ -90,4 +97,4 @@ const makePass = (getOutputs, ready, setSize, execute) => ({ const makePipeline = (context, steps) => steps.filter((f) => f != null).reduce((pipeline, f, i) => [...pipeline, f(context, i == 0 ? null : pipeline[i - 1].getOutputs)], []); -export { getCanvasSize, makePassFBO, make1DTexture, loadTexture, loadShader, makeUniformBuffer, makePass, makePipeline, makeBindGroup }; +export { getCanvasSize, makeRenderTarget, makeComputeTarget, make1DTexture, loadTexture, loadShader, makeUniformBuffer, makePass, makePipeline, makeBindGroup }; diff --git a/shaders/wgsl/postProcessingPass.wgsl b/shaders/wgsl/postProcessingPass.wgsl new file mode 100644 index 0000000..b32ffd9 --- /dev/null +++ b/shaders/wgsl/postProcessingPass.wgsl @@ -0,0 +1,66 @@ +[[block]] struct Config { + foo : i32; +}; + +// The properties that change over time get their own buffer. +[[block]] struct Time { + seconds : f32; + frames : i32; +}; + +[[group(0), binding(0)]] var config : Config; +[[group(0), binding(1)]] var time : Time; + +[[group(0), binding(2)]] var inputTex : texture_2d; +[[group(0), binding(3)]] var outputTex : texture_storage_2d; + +// Shader params + +struct ComputeInput { + [[builtin(global_invocation_id)]] id : vec3; +}; + +// Constants + +let NUM_VERTICES_PER_QUAD : i32 = 6; // 2 * 3 +let PI : f32 = 3.14159265359; +let TWO_PI : f32 = 6.28318530718; +let SQRT_2 : f32 = 1.4142135623730951; +let SQRT_5 : f32 = 2.23606797749979; + +// Helper functions for generating randomness, borrowed from elsewhere + +fn randomFloat( uv : vec2 ) -> f32 { + let a = 12.9898; + let b = 78.233; + let c = 43758.5453; + let dt = dot( uv, vec2( a, b ) ); + let sn = dt % PI; + return fract(sin(sn) * c); +} + +fn randomVec2( uv : vec2 ) -> vec2 { + return fract(vec2(sin(uv.x * 591.32 + uv.y * 154.077), cos(uv.x * 391.32 + uv.y * 49.077))); +} + +fn wobble(x : f32) -> f32 { + return x + 0.3 * sin(SQRT_2 * x) + 0.2 * sin(SQRT_5 * x); +} + +[[stage(compute), workgroup_size(32, 1, 1)]] fn computeMain(input : ComputeInput) { + + // Resolve the invocation ID to a single cell + var coord = vec2(input.id.xy); + var screenSize = textureDimensions(inputTex); + + if (coord.x >= screenSize.x) { + return; + } + + var foo = config.foo; + var seconds = time.seconds; + + var inputColor = textureLoad(inputTex, coord, 0); + var outputColor = inputColor; + textureStore(outputTex, coord, outputColor); +}