From 1c1b1e4f039ba129442b5eaddc4ef798b8fd7aa5 Mon Sep 17 00:00:00 2001 From: Rezmason Date: Mon, 1 Nov 2021 00:28:40 -0700 Subject: [PATCH] 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 | 6 -- js/webgpu/main.js | 37 ++++++++----- .../{rainRenderPass.wgsl => rainPass.wgsl} | 55 ++++++++++++++----- 3 files changed, 63 insertions(+), 35 deletions(-) rename shaders/wgsl/{rainRenderPass.wgsl => rainPass.wgsl} (87%) diff --git a/TODO.txt b/TODO.txt index cd3d9d2..8ade7f2 100644 --- a/TODO.txt +++ b/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 diff --git a/js/webgpu/main.js b/js/webgpu/main.js index e9e0118..6b98631 100644 --- a/js/webgpu/main.js +++ b/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", "mat4x4", "mat4x4"]); const sceneBuffer = makeUniformBuffer(device, sceneLayout); + const cellBufferDescriptor = { + size: numCells * std140(["vec4"]).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({ - layout: rainRenderPipeline.getBindGroupLayout(0), - entries: [configBuffer, timeBuffer, sceneBuffer, msdfSampler, msdfTexture.createView()] + const computeBindGroup = device.createBindGroup({ + layout: rainComputePipeline.getBindGroupLayout(0), + 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({ - layout: rainComputePipeline.getBindGroupLayout(0), - entries: [configBuffer, timeBuffer] + const renderBindGroup = device.createBindGroup({ + layout: rainRenderPipeline.getBindGroupLayout(0), + 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(); diff --git a/shaders/wgsl/rainRenderPass.wgsl b/shaders/wgsl/rainPass.wgsl similarity index 87% rename from shaders/wgsl/rainRenderPass.wgsl rename to shaders/wgsl/rainPass.wgsl index f14a9f6..02ae58c 100644 --- a/shaders/wgsl/rainRenderPass.wgsl +++ b/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; volumetric : i32; }; -[[group(0), binding(0)]] var config : Config; [[block]] struct Time { seconds : f32; frames : i32; }; -[[group(0), binding(1)]] var time : Time; [[block]] struct Scene { screenSize : vec2; camera : mat4x4; transform : mat4x4; }; -[[group(0), binding(2)]] var scene : Scene; +[[block]] struct CellData { + cells: array>; +}; + +// Shared bindings +[[group(0), binding(0)]] var config : Config; +[[group(0), binding(1)]] var time : Time; + +// Compute bindings +[[group(0), binding(2)]] var cellsPing_RW : CellData; +[[group(0), binding(3)]] var cellsPong_RW : CellData; + +// Render bindings +[[group(0), binding(2)]] var scene : Scene; [[group(0), binding(3)]] var msdfSampler : sampler; [[group(0), binding(4)]] var msdfTexture : texture_2d; - - +[[group(0), binding(5)]] var cellsPing_RO : CellData; +[[group(0), binding(6)]] var cellsPong_RO : CellData; // Shader params @@ -84,6 +87,14 @@ struct FragOutput { [[location(0)]] color : vec4; }; +// 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 { @@ -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((1.0 + time.seconds * 0.1) % 1.0, 0.0, 0.0, 0.0); + cellsPong_RW.cells[i] = vec4((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(1.0, 0.0, randomFloat(vec2(quadPosition.x, 1.0)), 0.0); // TODO : texture2D(state, quadPosition / quadGridSize); + var vGlyph: vec4; + 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 { [[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;