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

@@ -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 renderTarget = makeRenderTarget(device, 1, 1, canvasFormat);
-	const getOutputs = () => ({ ...getInputs(), bloom: fbo }); // TODO
+	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
+// A pyramid is just an array of Targets, where each Target is half the width
-// and half the height of the FBO below it.
+// 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 highPassTarget = highPassPyramid[i];
-				const hBlurFBO = hBlurPyramid[i];
+				const hBlurTarget = hBlurPyramid[i];
-				const vBlurFBO = vBlurPyramid[i];
+				const vBlurTarget = vBlurPyramid[i];
-				highPass({ fbo: highPassFBO, frag: highPassFrag.text(), tex: inputs.primary });
+				highPass({ fbo: highPassTarget, frag: highPassFrag.text(), tex: inputs.primary });
-				blur({ fbo: hBlurFBO, frag: blurFrag.text(), tex: highPassFBO, direction: [1, 0] });
+				blur({ fbo: hBlurTarget, frag: blurFrag.text(), tex: highPassTarget, direction: [1, 0] });
-				blur({ fbo: vBlurFBO, frag: blurFrag.text(), tex: hBlurFBO, direction: [0, 1] });
+				blur({ fbo: vBlurTarget, frag: blurFrag.text(), tex: hBlurTarget, direction: [0, 1] });
 			}
 
 			sumPyramid();

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;
 

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) => {

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));
 

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) => {

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);
+};

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);

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 = () => ({

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 = () => ({

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 };

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<uniform> config : Config;
+[[group(0), binding(1)]] var<uniform> time : Time;
+
+[[group(0), binding(2)]] var inputTex : texture_2d<f32>;
+[[group(0), binding(3)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
+
+// Shader params
+
+struct ComputeInput {
+	[[builtin(global_invocation_id)]] id : vec3<u32>;
+};
+
+// 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> ) -> f32 {
+	let a = 12.9898;
+	let b = 78.233;
+	let c = 43758.5453;
+	let dt = dot( uv, vec2<f32>( a, b ) );
+	let sn = dt % PI;
+	return fract(sin(sn) * c);
+}
+
+fn randomVec2( uv : vec2<f32> ) -> vec2<f32> {
+	return fract(vec2<f32>(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<i32>(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);
+}