We have ping-pong! And all the rain pass's shaders can coexist in one module! And I met and received ample help from @kainino0x and others on the WebGPU Matrix chat!

TODO.txt

@@ -3,12 +3,6 @@ TODO:
 WebGPU
   First decent rainRender
     Port compute pass 100%
-      Compute entry point can live in the same wgsl file, I think
-        ping-pong buffers
-      Bind both ping-pong buffers in render pass
-        The frame integer can be used by the render code to reference the correct one
-    Reorder the config fields
-    Reconsider how the bind groups are organized?
     Render target
       Resize accordingly
     Put in its own module

js/webgpu/main.js

@@ -31,17 +31,18 @@ export default async (canvas, config) => {
 	canvasContext.configure(canvasConfig);
 
 	const msdfTexturePromise = loadTexture(device, config.glyphTexURL);
-	const rainRenderShaderPromise = fetch("shaders/wgsl/rainRenderPass.wgsl").then((response) => response.text());
+	const rainShaderPromise = fetch("shaders/wgsl/rainPass.wgsl").then((response) => response.text());
 
 	// The volumetric mode multiplies the number of columns
 	// to reach the desired density, and then overlaps them
 	const volumetric = config.volumetric;
 	const density = volumetric && config.effect !== "none" ? config.density : 1;
 	const gridSize = [config.numColumns * density, config.numColumns];
+	const numCells = gridSize[0] * gridSize[1];
 
 	// The volumetric mode requires us to create a grid of quads,
 	// rather than a single quad for our geometry
-	const numQuads = volumetric ? gridSize[0] * gridSize[1] : 1;
+	const numQuads = volumetric ? numCells : 1;
 
 	// Various effect-related values
 	const rippleType = config.rippleTypeName in rippleTypes ? rippleTypes[config.rippleTypeName] : -1;
@@ -102,6 +103,13 @@ export default async (canvas, config) => {
 	const sceneLayout = std140(["vec2<f32>", "mat4x4<f32>", "mat4x4<f32>"]);
 	const sceneBuffer = makeUniformBuffer(device, sceneLayout);
 
+	const cellBufferDescriptor = {
+		size: numCells * std140(["vec4<f32>"]).size, // TODO: Is this correct?
+		usage: GPUBufferUsage.STORAGE,
+	};
+	const cellsPingBuffer = device.createBuffer(cellBufferDescriptor);
+	const cellsPongBuffer = device.createBuffer(cellBufferDescriptor);
+
 	const transform = mat4.create();
 	mat4.translate(transform, transform, vec3.fromValues(0, 0, -1));
 	const camera = mat4.create();
@@ -120,13 +128,13 @@ export default async (canvas, config) => {
 		minFilter: "linear",
 	});
 
-	const [msdfTexture, rainRenderShader] = await Promise.all([msdfTexturePromise, rainRenderShaderPromise]);
+	const [msdfTexture, rainShader] = await Promise.all([msdfTexturePromise, rainShaderPromise]);
 
-	const rainRenderShaderModule = device.createShaderModule({ code: rainRenderShader });
+	const rainShaderModule = device.createShaderModule({ code: rainShader });
 
 	const rainComputePipeline = device.createComputePipeline({
 		compute: {
-			module: rainRenderShaderModule,
+			module: rainShaderModule,
 			entryPoint: "computeMain",
 		},
 	});
@@ -139,11 +147,11 @@ export default async (canvas, config) => {
 
 	const rainRenderPipeline = device.createRenderPipeline({
 		vertex: {
-			module: rainRenderShaderModule,
+			module: rainShaderModule,
 			entryPoint: "vertMain",
 		},
 		fragment: {
-			module: rainRenderShaderModule,
+			module: rainShaderModule,
 			entryPoint: "fragMain",
 			targets: [
 				{
@@ -157,9 +165,9 @@ export default async (canvas, config) => {
 		},
 	});
 
-	const renderBindGroup = device.createBindGroup({
+	const computeBindGroup = device.createBindGroup({
-		layout: rainRenderPipeline.getBindGroupLayout(0),
+		layout: rainComputePipeline.getBindGroupLayout(0),
-		entries: [configBuffer, timeBuffer, sceneBuffer, msdfSampler, msdfTexture.createView()]
+		entries: [configBuffer, timeBuffer, cellsPingBuffer, cellsPongBuffer]
 			.map((resource) => (resource instanceof GPUBuffer ? { buffer: resource } : resource))
 			.map((resource, binding) => ({
 				binding,
@@ -167,9 +175,9 @@ export default async (canvas, config) => {
 			})),
 	});
 
-	const computeBindGroup = device.createBindGroup({
+	const renderBindGroup = device.createBindGroup({
-		layout: rainComputePipeline.getBindGroupLayout(0),
+		layout: rainRenderPipeline.getBindGroupLayout(0),
-		entries: [configBuffer, timeBuffer]
+		entries: [configBuffer, timeBuffer, sceneBuffer, msdfSampler, msdfTexture.createView(), cellsPingBuffer, cellsPongBuffer]
 			.map((resource) => (resource instanceof GPUBuffer ? { buffer: resource } : resource))
 			.map((resource, binding) => ({
 				binding,
@@ -212,8 +220,6 @@ export default async (canvas, config) => {
 		device.queue.writeBuffer(timeBuffer, 0, timeLayout.build([now / 1000, frame]));
 		frame++;
 
-		renderPassConfig.colorAttachments[0].view = canvasContext.getCurrentTexture().createView();
-
 		const encoder = device.createCommandEncoder();
 
 		const computePass = encoder.beginComputePass();
@@ -222,6 +228,7 @@ export default async (canvas, config) => {
 		computePass.dispatch(...gridSize, 1);
 		computePass.endPass();
 
+		renderPassConfig.colorAttachments[0].view = canvasContext.getCurrentTexture().createView();
 		const renderPass = encoder.beginRenderPass(renderPassConfig);
 		renderPass.executeBundles(renderBundles);
 		renderPass.endPass();

shaders/wgsl/rainPass.wgsl

@@ -1,11 +1,3 @@
-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;
-
-// Bound resources
-
 [[block]] struct Config {
 	// common
 	animationSpeed : f32;
@@ -43,25 +35,36 @@ let SQRT_5 : f32 = 2.23606797749979;
 	slantVec : vec2<f32>;
 	volumetric : i32;
 };
-[[group(0), binding(0)]] var<uniform> config : Config;
 
 [[block]] struct Time {
 	seconds : f32;
 	frames : i32;
 };
-[[group(0), binding(1)]] var<uniform> time : Time;
 
 [[block]] struct Scene {
 	screenSize : vec2<f32>;
 	camera : mat4x4<f32>;
 	transform : mat4x4<f32>;
 };
-[[group(0), binding(2)]] var<uniform> scene : Scene;
 
+[[block]] struct CellData {
+	cells: array<vec4<f32>>;
+};
+
+// Shared bindings
+[[group(0), binding(0)]] var<uniform> config : Config;
+[[group(0), binding(1)]] var<uniform> time : Time;
+
+// Compute bindings
+[[group(0), binding(2)]] var<storage, read_write> cellsPing_RW : CellData;
+[[group(0), binding(3)]] var<storage, read_write> cellsPong_RW : CellData;
+
+// Render bindings
+[[group(0), binding(2)]] var<uniform> scene : Scene;
 [[group(0), binding(3)]] var msdfSampler : sampler;
 [[group(0), binding(4)]] var msdfTexture : texture_2d<f32>;
-
+[[group(0), binding(5)]] var<storage, read> cellsPing_RO : CellData;
-
+[[group(0), binding(6)]] var<storage, read> cellsPong_RO : CellData;
 
 // Shader params
 
@@ -84,6 +87,14 @@ struct FragOutput {
 	[[location(0)]] color : vec4<f32>;
 };
 
+// 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 {
@@ -106,8 +117,16 @@ fn wobble(x : f32) -> f32 {
 // Compute shader
 
 [[stage(compute), workgroup_size(1, 1, 1)]] fn computeMain(input : ComputeInput) {
+	var animationSpeed = config.animationSpeed; // TODO: remove
 	var hasSun = bool(config.hasSun); // TODO: remove
 	var seconds = time.seconds; // TODO: remove
+
+	var row = i32(input.id.y);
+	var column = i32(input.id.x);
+	var i = row * i32(config.gridSize.x);
+
+	cellsPing_RW.cells[i] = vec4<f32>((1.0 + time.seconds * 0.1) % 1.0, 0.0, 0.0, 0.0);
+	cellsPong_RW.cells[i] = vec4<f32>((0.5 + time.seconds * 0.1) % 1.0, 0.5, 0.0, 0.0);
 }
 
 // Vertex shader
@@ -139,7 +158,12 @@ fn wobble(x : f32) -> f32 {
 	var uv = (quadPosition + quadCorner) / quadGridSize;
 
 	// Retrieve the quad's glyph data
-	var vGlyph = vec4<f32>(1.0, 0.0, randomFloat(vec2<f32>(quadPosition.x, 1.0)), 0.0); // TODO : texture2D(state, quadPosition / quadGridSize);
+	var vGlyph: vec4<f32>;
+	if ((time.frames / 100) % 2 == 0) {
+		vGlyph = cellsPing_RO.cells[quadIndex];
+	} else {
+		vGlyph = cellsPong_RO.cells[quadIndex];
+	}
 
 	// Calculate the quad's depth
 	var quadDepth = 0.0;
@@ -197,6 +221,9 @@ fn getSymbolUV(glyphCycle : f32) -> vec2<f32> {
 
 [[stage(fragment)]] fn fragMain(input : VertOutput) -> FragOutput {
 
+	var firstCellA = cellsPing_RO.cells[0]; // TODO: remove
+	var firstCellB = cellsPong_RO.cells[0]; // TODO: remove
+
 	var volumetric = bool(config.volumetric);
 	var uv = input.uv;