All the post processing passes are now based on compute pipelines instead of render pipelines.

TODO.txt

@@ -1,7 +1,6 @@
 TODO:
 
 WebGPU
-  Switch post processing to compute shaders
   blur pass
   Update links in issues
   Get rid of end pass once it's possible to copy a bgra8unorm to a canvas texture

js/webgpu/bloomPass.js

@@ -1,5 +1,5 @@
 import { structs } from "/lib/gpu-buffer.js";
-import { loadShader, makeUniformBuffer, makeBindGroup, makeRenderTarget, makePass } from "./utils.js";
+import { loadShader, makeUniformBuffer, makeBindGroup, makeComputeTarget, 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.
@@ -11,13 +11,13 @@ const levelStrengths = Array(pyramidHeight)
 	.reverse();
 
 export default (context, getInputs) => {
-	const { config, device, canvasFormat } = context;
+	const { config, device } = context;
 
-	const enabled = config.bloomSize > 0 && config.bloomStrength > 0;
+	const enabled = false; // config.bloomSize > 0 && config.bloomStrength > 0;
 
 	// If there's no bloom to apply, return a no-op pass with an empty bloom texture
 	if (!enabled) {
-		const emptyTexture = makeRenderTarget(device, 1, 1, canvasFormat);
+		const emptyTexture = makeComputeTarget(device, 1, 1);
 		const getOutputs = () => ({ ...getInputs(), bloom: emptyTexture });
 		return makePass(getOutputs);
 	}
@@ -26,11 +26,11 @@ export default (context, getInputs) => {
 
 	// TODO: generate sum shader code
 
-	const renderTarget = makeRenderTarget(device, 1, 1, canvasFormat);
-	const getOutputs = () => ({ ...getInputs(), bloom: renderTarget }); // TODO
+	const computeTarget = makeComputeTarget(device, 1, 1);
+	const getOutputs = () => ({ ...getInputs(), bloom: computeTarget }); // TODO
 
-	let blurRenderPipeline;
-	let sumRenderPipeline;
+	let blurPipeline;
+	let sumPipeline;
 
 	const ready = (async () => {
 		const [blurShader] = await Promise.all(assets);
@@ -74,7 +74,7 @@ export default (context, getInputs) => {
 const makePyramid = (regl, height, halfFloat) =>
 	Array(height)
 		.fill()
-		.map((_) => makeRenderTarget(regl, halfFloat));
+		.map((_) => makePassFBO(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)));
@@ -85,7 +85,7 @@ export default ({ regl, config }, inputs) => {
 	const highPassPyramid = makePyramid(regl, pyramidHeight, config.useHalfFloat);
 	const hBlurPyramid = makePyramid(regl, pyramidHeight, config.useHalfFloat);
 	const vBlurPyramid = makePyramid(regl, pyramidHeight, config.useHalfFloat);
-	const output = makeRenderTarget(regl, config.useHalfFloat);
+	const output = makePassFBO(regl, config.useHalfFloat);
 
 	// The high pass restricts the blur to bright things in our input texture.
 	const highPassFrag = loadText("shaders/glsl/highPass.frag.glsl");

js/webgpu/imagePass.js

@@ -1,35 +1,25 @@
-import { loadTexture, loadShader, makeBindGroup, makeRenderTarget, makePass } from "./utils.js";
+import { makeComputeTarget, loadTexture, loadShader, makeUniformBuffer, makeBindGroup, makePass } from "./utils.js";
 
 // Multiplies the rendered rain and bloom by a loaded in image
 
 const defaultBGURL = "https://upload.wikimedia.org/wikipedia/commons/thumb/0/0a/Flammarion_Colored.jpg/917px-Flammarion_Colored.jpg";
-const numVerticesPerQuad = 2 * 3;
 
 export default (context, getInputs) => {
-	const { config, device, canvasFormat } = context;
+	const { config, device } = context;
+
+	const bgURL = "bgURL" in config ? config.bgURL : defaultBGURL;
+	const assets = [loadTexture(device, bgURL), loadShader(device, "shaders/wgsl/imagePass.wgsl")];
 
 	const linearSampler = device.createSampler({
 		magFilter: "linear",
 		minFilter: "linear",
 	});
 
-	const renderPassConfig = {
-		colorAttachments: [
-			{
-				view: null,
-				loadValue: { r: 0, g: 0, b: 0, a: 1 },
-				storeOp: "store",
-			},
-		],
-	};
-
-	let renderPipeline;
+	let computePipeline;
 	let output;
+	let screenSize;
 	let backgroundTex;
 
-	const bgURL = "bgURL" in config ? config.bgURL : defaultBGURL;
-	const assets = [loadTexture(device, bgURL), loadShader(device, "shaders/wgsl/imagePass.wgsl")];
-
 	const getOutputs = () => ({
 		primary: output,
 	});
@@ -39,39 +29,36 @@ export default (context, getInputs) => {
 
 		backgroundTex = bgTex;
 
-		renderPipeline = device.createRenderPipeline({
-			vertex: {
+		computePipeline = device.createComputePipeline({
+			compute: {
 				module: imageShader.module,
-				entryPoint: "vertMain",
-			},
-			fragment: {
-				module: imageShader.module,
-				entryPoint: "fragMain",
-				targets: [
-					{
-						format: canvasFormat,
-					},
-				],
+				entryPoint: "computeMain",
 			},
 		});
 	})();
 
 	const setSize = (width, height) => {
 		output?.destroy();
-		output = makeRenderTarget(device, width, height, canvasFormat);
+		output = makeComputeTarget(device, width, height);
+		screenSize = [width, height];
 	};
 
 	const execute = (encoder) => {
 		const inputs = getInputs();
 		const tex = inputs.primary;
-		const bloomTex = inputs.bloom; // TODO: bloom
-		const renderBindGroup = makeBindGroup(device, renderPipeline, 0, [linearSampler, tex.createView(), bloomTex.createView(), backgroundTex.createView()]);
-		renderPassConfig.colorAttachments[0].view = output.createView();
-		const renderPass = encoder.beginRenderPass(renderPassConfig);
-		renderPass.setPipeline(renderPipeline);
-		renderPass.setBindGroup(0, renderBindGroup);
-		renderPass.draw(numVerticesPerQuad, 1, 0, 0);
-		renderPass.endPass();
+		const bloomTex = inputs.bloom;
+		const computePass = encoder.beginComputePass();
+		computePass.setPipeline(computePipeline);
+		const computeBindGroup = makeBindGroup(device, computePipeline, 0, [
+			linearSampler,
+			tex.createView(),
+			bloomTex.createView(),
+			backgroundTex.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/main.js

@@ -7,7 +7,6 @@ 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 = {
@@ -53,7 +52,7 @@ export default async (canvas, config) => {
 	};
 
 	const effectName = config.effect in effects ? config.effect : "plain";
-	const pipeline = makePipeline(context, [makeRain, makeBloomPass, effects[effectName], makePostProcessingPass, makeEndPass]);
+	const pipeline = makePipeline(context, [makeRain, makeBloomPass, effects[effectName], 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, makeRenderTarget, makePass } from "./utils.js";
+import { loadShader, makeUniformBuffer, makeBindGroup, makeComputeTarget, 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
@@ -15,8 +15,6 @@ const colorToRGB = ([hue, saturation, lightness]) => {
 	return [f(0), f(8), f(4)];
 };
 
-const numVerticesPerQuad = 2 * 3;
-
 const makePalette = (device, paletteUniforms, entries) => {
 	const PALETTE_SIZE = 512;
 	const paletteColors = Array(PALETTE_SIZE);
@@ -78,27 +76,18 @@ const makePalette = (device, paletteUniforms, entries) => {
 // in screen space.
 
 export default (context, getInputs) => {
-	const { config, device, timeBuffer, canvasFormat } = context;
+	const { config, device, timeBuffer } = context;
 
 	const linearSampler = device.createSampler({
 		magFilter: "linear",
 		minFilter: "linear",
 	});
 
-	const renderPassConfig = {
-		colorAttachments: [
-			{
-				view: null,
-				loadValue: { r: 0, g: 0, b: 0, a: 1 },
-				storeOp: "store",
-			},
-		],
-	};
-
-	let renderPipeline;
+	let computePipeline;
 	let configBuffer;
 	let paletteBuffer;
 	let output;
+	let screenSize;
 
 	const getOutputs = () => ({
 		primary: output,
@@ -109,19 +98,10 @@ export default (context, getInputs) => {
 	const ready = (async () => {
 		const [paletteShader] = await Promise.all(assets);
 
-		renderPipeline = device.createRenderPipeline({
-			vertex: {
+		computePipeline = device.createComputePipeline({
+			compute: {
 				module: paletteShader.module,
-				entryPoint: "vertMain",
-			},
-			fragment: {
-				module: paletteShader.module,
-				entryPoint: "fragMain",
-				targets: [
-					{
-						format: canvasFormat,
-					},
-				],
+				entryPoint: "computeMain",
 			},
 		});
 
@@ -135,28 +115,28 @@ export default (context, getInputs) => {
 
 	const setSize = (width, height) => {
 		output?.destroy();
-		output = makeRenderTarget(device, width, height, canvasFormat);
+		output = makeComputeTarget(device, width, height);
+		screenSize = [width, height];
 	};
 
 	const execute = (encoder) => {
 		const inputs = getInputs();
 		const tex = inputs.primary;
-		const bloomTex = inputs.bloom; // TODO: bloom
-		const renderBindGroup = makeBindGroup(device, renderPipeline, 0, [
+		const bloomTex = inputs.bloom;
+		const computePass = encoder.beginComputePass();
+		computePass.setPipeline(computePipeline);
+		const computeBindGroup = makeBindGroup(device, computePipeline, 0, [
 			configBuffer,
 			paletteBuffer,
 			timeBuffer,
 			linearSampler,
 			tex.createView(),
 			bloomTex.createView(),
+			output.createView(),
 		]);
-
-		renderPassConfig.colorAttachments[0].view = output.createView();
-		const renderPass = encoder.beginRenderPass(renderPassConfig);
-		renderPass.setPipeline(renderPipeline);
-		renderPass.setBindGroup(0, renderBindGroup);
-		renderPass.draw(numVerticesPerQuad, 1, 0, 0);
-		renderPass.endPass();
+		computePass.setBindGroup(0, computeBindGroup);
+		computePass.dispatch(Math.ceil(screenSize[0] / 32), screenSize[1], 1);
+		computePass.endPass();
 	};
 
 	return makePass(getOutputs, ready, setSize, execute);

js/webgpu/postProcessingPass.js

@@ -1,52 +0,0 @@
-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/resurrectionPass.js

@@ -1,5 +1,5 @@
 import { structs } from "/lib/gpu-buffer.js";
-import { loadShader, makeUniformBuffer, makeRenderTarget, makePass } from "./utils.js";
+import { loadShader, makeUniformBuffer, makeComputeTarget, makeBindGroup, makePass } from "./utils.js";
 
 // Matrix Resurrections isn't in theaters yet,
 // and this version of the effect is still a WIP.
@@ -12,45 +12,27 @@ import { loadShader, makeUniformBuffer, makeRenderTarget, makePass } from "./uti
 const numVerticesPerQuad = 2 * 3;
 
 export default (context, getInputs) => {
-	const { config, device, timeBuffer, canvasFormat } = context;
+	const { config, device, timeBuffer } = context;
 
 	const linearSampler = device.createSampler({
 		magFilter: "linear",
 		minFilter: "linear",
 	});
 
-	const renderPassConfig = {
-		colorAttachments: [
-			{
-				view: null,
-				loadValue: { r: 0, g: 0, b: 0, a: 1 },
-				storeOp: "store",
-			},
-		],
-	};
-
-	let renderPipeline;
+	let computePipeline;
 	let configBuffer;
 	let output;
+	let screenSize;
 
 	const assets = [loadShader(device, "shaders/wgsl/resurrectionPass.wgsl")];
 
 	const ready = (async () => {
 		const [resurrectionShader] = await Promise.all(assets);
 
-		renderPipeline = device.createRenderPipeline({
-			vertex: {
+		computePipeline = device.createComputePipeline({
+			compute: {
 				module: resurrectionShader.module,
-				entryPoint: "vertMain",
-			},
-			fragment: {
-				module: resurrectionShader.module,
-				entryPoint: "fragMain",
-				targets: [
-					{
-						format: canvasFormat,
-					},
-				],
+				entryPoint: "computeMain",
 			},
 		});
 
@@ -60,7 +42,8 @@ export default (context, getInputs) => {
 
 	const setSize = (width, height) => {
 		output?.destroy();
-		output = makeRenderTarget(device, width, height, canvasFormat);
+		output = makeComputeTarget(device, width, height);
+		screenSize = [width, height];
 	};
 
 	const getOutputs = () => ({
@@ -70,15 +53,20 @@ export default (context, getInputs) => {
 	const execute = (encoder) => {
 		const inputs = getInputs();
 		const tex = inputs.primary;
-		const bloomTex = inputs.bloom; // TODO: bloom
-		const renderBindGroup = makeBindGroup(device, renderPipeline, 0, [configBuffer, timeBuffer, linearSampler, tex.createView(), bloomTex.createView()]);
-
-		renderPassConfig.colorAttachments[0].view = output.createView();
-		const renderPass = encoder.beginRenderPass(renderPassConfig);
-		renderPass.setPipeline(renderPipeline);
-		renderPass.setBindGroup(0, renderBindGroup);
-		renderPass.draw(numVerticesPerQuad, 1, 0, 0);
-		renderPass.endPass();
+		const bloomTex = inputs.bloom;
+		const computePass = encoder.beginComputePass();
+		computePass.setPipeline(computePipeline);
+		const computeBindGroup = makeBindGroup(device, computePipeline, 0, [
+			configBuffer,
+			timeBuffer,
+			linearSampler,
+			tex.createView(),
+			bloomTex.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/stripePass.js

@@ -1,5 +1,5 @@
 import { structs } from "/lib/gpu-buffer.js";
-import { loadShader, make1DTexture, makeUniformBuffer, makeBindGroup, makeRenderTarget, makePass } from "./utils.js";
+import { loadShader, make1DTexture, makeUniformBuffer, makeBindGroup, makeComputeTarget, makePass } from "./utils.js";
 
 // Multiplies the rendered rain and bloom by a 1D gradient texture
 // generated from the passed-in color sequence
@@ -38,7 +38,7 @@ const numVerticesPerQuad = 2 * 3;
 // in screen space.
 
 export default (context, getInputs) => {
-	const { config, device, timeBuffer, canvasFormat } = context;
+	const { config, device, timeBuffer } = context;
 
 	// Expand and convert stripe colors into 1D texture data
 	const input =
@@ -55,38 +55,20 @@ export default (context, getInputs) => {
 		minFilter: "linear",
 	});
 
-	const renderPassConfig = {
-		colorAttachments: [
-			{
-				view: null,
-				loadValue: { r: 0, g: 0, b: 0, a: 1 },
-				storeOp: "store",
-			},
-		],
-	};
-
-	let renderPipeline;
+	let computePipeline;
 	let configBuffer;
 	let output;
+	let screenSize;
 
 	const assets = [loadShader(device, "shaders/wgsl/stripePass.wgsl")];
 
 	const ready = (async () => {
 		const [stripeShader] = await Promise.all(assets);
 
-		renderPipeline = device.createRenderPipeline({
-			vertex: {
+		computePipeline = device.createComputePipeline({
+			compute: {
 				module: stripeShader.module,
-				entryPoint: "vertMain",
-			},
-			fragment: {
-				module: stripeShader.module,
-				entryPoint: "fragMain",
-				targets: [
-					{
-						format: canvasFormat,
-					},
-				],
+				entryPoint: "computeMain",
 			},
 		});
 
@@ -96,7 +78,8 @@ export default (context, getInputs) => {
 
 	const setSize = (width, height) => {
 		output?.destroy();
-		output = makeRenderTarget(device, width, height, canvasFormat);
+		output = makeComputeTarget(device, width, height);
+		screenSize = [width, height];
 	};
 
 	const getOutputs = () => ({
@@ -106,22 +89,21 @@ export default (context, getInputs) => {
 	const execute = (encoder) => {
 		const inputs = getInputs();
 		const tex = inputs.primary;
-		const bloomTex = inputs.bloom; // TODO: bloom
-		const renderBindGroup = makeBindGroup(device, renderPipeline, 0, [
+		const bloomTex = inputs.bloom;
+		const computePass = encoder.beginComputePass();
+		computePass.setPipeline(computePipeline);
+		const computeBindGroup = makeBindGroup(device, computePipeline, 0, [
 			configBuffer,
 			timeBuffer,
 			linearSampler,
 			tex.createView(),
 			bloomTex.createView(),
 			stripeTexture.createView(),
+			output.createView(),
 		]);
-
-		renderPassConfig.colorAttachments[0].view = output.createView();
-		const renderPass = encoder.beginRenderPass(renderPassConfig);
-		renderPass.setPipeline(renderPipeline);
-		renderPass.setBindGroup(0, renderBindGroup);
-		renderPass.draw(numVerticesPerQuad, 1, 0, 0);
-		renderPass.endPass();
+		computePass.setBindGroup(0, computeBindGroup);
+		computePass.dispatch(Math.ceil(screenSize[0] / 32), screenSize[1], 1);
+		computePass.endPass();
 	};
 
 	return makePass(getOutputs, ready, setSize, execute);

shaders/wgsl/imagePass.wgsl

@@ -2,29 +2,30 @@
 [[group(0), binding(1)]] var tex : texture_2d<f32>;
 [[group(0), binding(2)]] var bloomTex : texture_2d<f32>;
 [[group(0), binding(3)]] var backgroundTex : texture_2d<f32>;
+[[group(0), binding(4)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
 
-struct VertOutput {
-	[[builtin(position)]] Position : vec4<f32>;
-	[[location(0)]] uv : vec2<f32>;
+struct ComputeInput {
+	[[builtin(global_invocation_id)]] id : vec3<u32>;
 };
 
-[[stage(vertex)]] fn vertMain([[builtin(vertex_index)]] index : u32) -> VertOutput {
-	var uv = vec2<f32>(f32(index % 2u), f32((index + 1u) % 6u / 3u));
-	var position = vec4<f32>(uv * 2.0 - 1.0, 1.0, 1.0);
-	return VertOutput(position, uv);
-}
+[[stage(compute), workgroup_size(32, 1, 1)]] fn computeMain(input : ComputeInput) {
 
-[[stage(fragment)]] fn fragMain(input : VertOutput) -> [[location(0)]] vec4<f32> {
+	// Resolve the invocation ID to a single cell
+	var coord = vec2<i32>(input.id.xy);
+	var screenSize = textureDimensions(tex);
 
-	var uv = input.uv;
-	uv.y = 1.0 - uv.y;
+	if (coord.x >= screenSize.x) {
+		return;
+	}
 
-	var bgColor = textureSample( backgroundTex, linearSampler, uv ).rgb;
+	var uv = vec2<f32>(coord) / vec2<f32>(screenSize);
+
+	var bgColor = textureSampleLevel( backgroundTex, linearSampler, uv, 0.0 ).rgb;
 
 	// Combine the texture and bloom, then blow it out to reveal more of the image
-	var brightness = min(1.0, textureSample( tex, linearSampler, uv ).r * 2.0);
-	brightness = brightness + textureSample( bloomTex, linearSampler, uv ).r;
+	var brightness = min(1.0, textureSampleLevel( tex, linearSampler, uv, 0.0 ).r * 2.0);
+	brightness = brightness + textureSampleLevel( bloomTex, linearSampler, uv, 0.0 ).r;
 	brightness = pow(brightness, 1.5);
 
-	return vec4<f32>(bgColor * brightness, 1.0);
+	textureStore(outputTex, coord, vec4<f32>(bgColor * brightness, 1.0));
 }

shaders/wgsl/palettePass.wgsl

@@ -18,10 +18,10 @@
 [[group(0), binding(3)]] var linearSampler : sampler;
 [[group(0), binding(4)]] var tex : texture_2d<f32>;
 [[group(0), binding(5)]] var bloomTex : texture_2d<f32>;
+[[group(0), binding(6)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
 
-struct VertOutput {
-	[[builtin(position)]] Position : vec4<f32>;
-	[[location(0)]] uv : vec2<f32>;
+struct ComputeInput {
+	[[builtin(global_invocation_id)]] id : vec3<u32>;
 };
 
 let PI : f32 = 3.14159265359;
@@ -35,18 +35,20 @@ fn randomFloat( uv : vec2<f32> ) -> f32 {
 	return fract(sin(sn) * c);
 }
 
-[[stage(vertex)]] fn vertMain([[builtin(vertex_index)]] index : u32) -> VertOutput {
-	var uv = vec2<f32>(f32(index % 2u), f32((index + 1u) % 6u / 3u));
-	var position = vec4<f32>(uv * 2.0 - 1.0, 1.0, 1.0);
-	return VertOutput(position, uv);
-}
 
-[[stage(fragment)]] fn fragMain(input : VertOutput) -> [[location(0)]] vec4<f32> {
+[[stage(compute), workgroup_size(32, 1, 1)]] fn computeMain(input : ComputeInput) {
 
-	var uv = input.uv;
-	uv.y = 1.0 - uv.y;
+	// Resolve the invocation ID to a single cell
+	var coord = vec2<i32>(input.id.xy);
+	var screenSize = textureDimensions(tex);
 
-	var brightnessRGB = textureSample( tex, linearSampler, uv ) + textureSample( bloomTex, linearSampler, uv );
+	if (coord.x >= screenSize.x) {
+		return;
+	}
+
+	var uv = vec2<f32>(coord) / vec2<f32>(screenSize);
+
+	var brightnessRGB = textureSampleLevel( tex, linearSampler, uv, 0.0 ) + textureSampleLevel( bloomTex, linearSampler, uv, 0.0 );
 
 	// Combine the texture and bloom
 	var brightness = brightnessRGB.r + brightnessRGB.g + brightnessRGB.b;
@@ -57,5 +59,6 @@ fn randomFloat( uv : vec2<f32> ) -> f32 {
 	var paletteIndex = clamp(i32(brightness * 512.0), 0, 511);
 
 	// Map the brightness to a position in the palette texture
-	return vec4<f32>(palette.colors[paletteIndex] + config.backgroundColor, 1.0);
+	textureStore(outputTex, coord, vec4<f32>(palette.colors[paletteIndex] + config.backgroundColor, 1.0));
 }
+

shaders/wgsl/postProcessingPass.wgsl

@@ -1,66 +0,0 @@
-[[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);
-}

shaders/wgsl/resurrectionPass.wgsl

@@ -13,10 +13,10 @@
 [[group(0), binding(2)]] var linearSampler : sampler;
 [[group(0), binding(3)]] var tex : texture_2d<f32>;
 [[group(0), binding(4)]] var bloomTex : texture_2d<f32>;
+[[group(0), binding(5)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
 
-struct VertOutput {
-	[[builtin(position)]] Position : vec4<f32>;
-	[[location(0)]] uv : vec2<f32>;
+struct ComputeInput {
+	[[builtin(global_invocation_id)]] id : vec3<u32>;
 };
 
 let PI : f32 = 3.14159265359;
@@ -55,36 +55,36 @@ fn hslToRgb(h : f32, s : f32, l : f32) -> vec3<f32> {
 	);
 }
 
-[[stage(vertex)]] fn vertMain([[builtin(vertex_index)]] index : u32) -> VertOutput {
-	var uv = vec2<f32>(f32(index % 2u), f32((index + 1u) % 6u / 3u));
-	var position = vec4<f32>(uv * 2.0 - 1.0, 1.0, 1.0);
-	return VertOutput(position, uv);
-}
+[[stage(compute), workgroup_size(32, 1, 1)]] fn computeMain(input : ComputeInput) {
 
-[[stage(fragment)]] fn fragMain(input : VertOutput) -> [[location(0)]] vec4<f32> {
+	// Resolve the invocation ID to a single cell
+	var coord = vec2<i32>(input.id.xy);
+	var screenSize = textureDimensions(tex);
 
-	var uv = input.uv;
-	uv.y = 1.0 - uv.y;
+	if (coord.x >= screenSize.x) {
+		return;
+	}
 
+	var uv = vec2<f32>(coord) / vec2<f32>(screenSize);
 
 	// Mix the texture and bloom based on distance from center,
 	// to approximate a lens blur
 	var brightness = mix(
-		textureSample( tex, linearSampler, uv ).rgb,
-		textureSample( bloomTex, linearSampler, uv ).rgb,
-		(0.7 - length(input.uv - 0.5))
+		textureSampleLevel( tex, linearSampler, uv, 0.0 ).rgb,
+		textureSampleLevel( bloomTex, linearSampler, uv, 0.0 ).rgb,
+		(0.7 - length(uv - 0.5))
 	) * 1.25;
 
 	// Dither: subtract a random value from the brightness
 	brightness = brightness - randomFloat( uv + vec2<f32>(time.seconds) ) * config.ditherMagnitude;
 
 	// Calculate a hue based on distance from center
-	var hue = 0.35 + (length(input.uv - vec2<f32>(0.5, 1.0)) * -0.4 + 0.2);
+	var hue = 0.35 + (length(uv - vec2<f32>(0.5, 1.0)) * -0.4 + 0.2);
 
 	// Convert HSL to RGB
 	var rgb = hslToRgb(hue, 0.8, max(0., brightness.r)) * vec3<f32>(0.8, 1.0, 0.7);
 
 	// Calculate a separate RGB for upward-flowing glyphs
 	var resurrectionRGB = hslToRgb(0.13, 1.0, max(0., brightness.g) * 0.9);
-	return vec4<f32>(rgb + resurrectionRGB + config.backgroundColor, 1.0);
+	textureStore(outputTex, coord, vec4<f32>(rgb + resurrectionRGB + config.backgroundColor, 1.0));
 }

shaders/wgsl/stripePass.wgsl

@@ -14,10 +14,10 @@
 [[group(0), binding(3)]] var tex : texture_2d<f32>;
 [[group(0), binding(4)]] var bloomTex : texture_2d<f32>;
 [[group(0), binding(5)]] var stripeTexture : texture_2d<f32>;
+[[group(0), binding(6)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
 
-struct VertOutput {
-	[[builtin(position)]] Position : vec4<f32>;
-	[[location(0)]] uv : vec2<f32>;
+struct ComputeInput {
+	[[builtin(global_invocation_id)]] id : vec3<u32>;
 };
 
 let PI : f32 = 3.14159265359;
@@ -31,25 +31,26 @@ fn randomFloat( uv : vec2<f32> ) -> f32 {
 	return fract(sin(sn) * c);
 }
 
-[[stage(vertex)]] fn vertMain([[builtin(vertex_index)]] index : u32) -> VertOutput {
-	var uv = vec2<f32>(f32(index % 2u), f32((index + 1u) % 6u / 3u));
-	var position = vec4<f32>(uv * 2.0 - 1.0, 1.0, 1.0);
-	return VertOutput(position, uv);
-}
+[[stage(compute), workgroup_size(32, 1, 1)]] fn computeMain(input : ComputeInput) {
 
-[[stage(fragment)]] fn fragMain(input : VertOutput) -> [[location(0)]] vec4<f32> {
+	// Resolve the invocation ID to a single cell
+	var coord = vec2<i32>(input.id.xy);
+	var screenSize = textureDimensions(tex);
 
-	var uv = input.uv;
-	uv.y = 1.0 - uv.y;
+	if (coord.x >= screenSize.x) {
+		return;
+	}
 
-	var color = textureSample( stripeTexture, linearSampler, uv ).rgb;
+	var uv = vec2<f32>(coord) / vec2<f32>(screenSize);
+
+	var color = textureSampleLevel( stripeTexture, linearSampler, uv, 0.0 ).rgb;
 
 	// Combine the texture and bloom
-	var brightness = min(1.0, textureSample( tex, linearSampler, uv ).r * 2.0);
-	brightness = brightness + textureSample( bloomTex, linearSampler, uv ).r;
+	var brightness = min(1.0, textureSampleLevel( tex, linearSampler, uv, 0.0 ).r * 2.0);
+	brightness = brightness + textureSampleLevel( bloomTex, linearSampler, uv, 0.0 ).r;
 
 	// Dither: subtract a random value from the brightness
 	brightness = brightness - randomFloat( uv + vec2<f32>(time.seconds) ) * config.ditherMagnitude;
 
-	return vec4<f32>(color * brightness + config.backgroundColor, 1.0);
+	textureStore(outputTex, coord, vec4<f32>(color * brightness + config.backgroundColor, 1.0));
 }