From 348b07d23780ad7b3839a8d4869db500230a0ade Mon Sep 17 00:00:00 2001 From: Rezmason Date: Mon, 1 Nov 2021 08:38:48 -0700 Subject: [PATCH] Actually the ping-pong setup isn't necessary for this compute shader, since it does everything in-place. Undid all the y-flips I could find, and they all canceled one another out, except for the glyph UVs in the fragment shader. --- TODO.txt | 9 ++----- js/webgpu/main.js | 14 +++++----- shaders/wgsl/rainPass.wgsl | 55 +++++++++++++------------------------- 3 files changed, 27 insertions(+), 51 deletions(-) diff --git a/TODO.txt b/TODO.txt index 6e28cf7..5a7fece 100644 --- a/TODO.txt +++ b/TODO.txt @@ -1,13 +1,8 @@ TODO: WebGPU - First decent rainRender - Fix the rain direction - Fix polar glyph direction - Maybe re-flip everything - Render target - Resize accordingly - Put in its own module + Render targets + Resize accordingly Blur: compute or render? The other passes should be a breeze Just add them to the render bundle diff --git a/js/webgpu/main.js b/js/webgpu/main.js index 6b98631..a9969a3 100644 --- a/js/webgpu/main.js +++ b/js/webgpu/main.js @@ -103,12 +103,10 @@ 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? + const cellsBuffer = device.createBuffer({ + size: numCells * std140(["vec4"]).size, 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)); @@ -167,7 +165,7 @@ export default async (canvas, config) => { const computeBindGroup = device.createBindGroup({ layout: rainComputePipeline.getBindGroupLayout(0), - entries: [configBuffer, timeBuffer, cellsPingBuffer, cellsPongBuffer] + entries: [configBuffer, timeBuffer, cellsBuffer] .map((resource) => (resource instanceof GPUBuffer ? { buffer: resource } : resource)) .map((resource, binding) => ({ binding, @@ -177,7 +175,7 @@ export default async (canvas, config) => { const renderBindGroup = device.createBindGroup({ layout: rainRenderPipeline.getBindGroupLayout(0), - entries: [configBuffer, timeBuffer, sceneBuffer, msdfSampler, msdfTexture.createView(), cellsPingBuffer, cellsPongBuffer] + entries: [configBuffer, timeBuffer, sceneBuffer, msdfSampler, msdfTexture.createView(), cellsBuffer] .map((resource) => (resource instanceof GPUBuffer ? { buffer: resource } : resource)) .map((resource, binding) => ({ binding, @@ -225,7 +223,7 @@ export default async (canvas, config) => { const computePass = encoder.beginComputePass(); computePass.setPipeline(rainComputePipeline); computePass.setBindGroup(0, computeBindGroup); - computePass.dispatch(...gridSize, 1); + computePass.dispatch(Math.ceil(gridSize[0] / 32), gridSize[1], 1); computePass.endPass(); renderPassConfig.colorAttachments[0].view = canvasContext.getCurrentTexture().createView(); diff --git a/shaders/wgsl/rainPass.wgsl b/shaders/wgsl/rainPass.wgsl index 3a9fd39..0b0ea51 100644 --- a/shaders/wgsl/rainPass.wgsl +++ b/shaders/wgsl/rainPass.wgsl @@ -56,15 +56,13 @@ [[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; +[[group(0), binding(2)]] var cells_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; +[[group(0), binding(5)]] var cells_RO : CellData; // Shader params @@ -263,30 +261,22 @@ fn computeResult (isFirstFrame : bool, previousResult : vec4, glyphPos : ve return result; } -[[stage(compute), workgroup_size(1, 1, 1)]] fn computeMain(input : ComputeInput) { +[[stage(compute), workgroup_size(32, 1, 1)]] fn computeMain(input : ComputeInput) { var row = i32(input.id.y); var column = i32(input.id.x); + + if (column > i32(config.gridSize.x)) { + return; + } + var i = row * i32(config.gridSize.x) + column; var isFirstFrame = time.frames == 0; var glyphPos = vec2(f32(column), f32(row)); var screenPos = glyphPos / config.gridSize; - var previousResult : vec4; - - if (time.frames % 2 == 0) { - previousResult = cellsPong_RW.cells[i]; - } else { - previousResult = cellsPing_RW.cells[i]; - } - - var result = computeResult(isFirstFrame, previousResult, glyphPos, screenPos); - - if (time.frames % 2 == 0) { - cellsPing_RW.cells[i] = result; - } else { - cellsPong_RW.cells[i] = result; - } + var previousResult = cells_RW.cells[i]; + cells_RW.cells[i] = computeResult(isFirstFrame, previousResult, glyphPos, screenPos); } // Vertex shader @@ -318,12 +308,7 @@ fn computeResult (isFirstFrame : bool, previousResult : vec4, glyphPos : ve var uv = (quadPosition + quadCorner) / quadGridSize; // Retrieve the quad's glyph data - var vGlyph: vec4; - if (time.frames % 2 == 0) { - vGlyph = cellsPing_RO.cells[quadIndex]; - } else { - vGlyph = cellsPong_RO.cells[quadIndex]; - } + var vGlyph = cells_RO.cells[quadIndex]; // Calculate the quad's depth var quadDepth = 0.0; @@ -337,7 +322,6 @@ fn computeResult (isFirstFrame : bool, previousResult : vec4, glyphPos : ve worldPosition = worldPosition + quadCorner * vec2(config.density, 1.0); worldPosition = worldPosition / quadGridSize; worldPosition = (worldPosition - 0.5) * 2.0; - worldPosition.y = -worldPosition.y; // "Resurrected" columns are in the green channel, // and are vertically flipped (along with their glyphs) @@ -388,16 +372,18 @@ fn getSymbolUV(glyphCycle : f32) -> vec2 { if (!volumetric) { if (bool(config.isPolar)) { // Curve space to make the letters appear to radiate from up above - uv = (uv - 0.5) * 0.5; - uv.y = uv.y + 0.5; + uv = uv - 0.5; + uv = uv * 0.5; + uv.y = uv.y - 0.5; var radius = length(uv); var angle = atan2(uv.y, uv.x) / (2.0 * PI) + 0.5; uv = vec2(fract(angle * 4.0 - 0.5), 1.5 * (1.0 - sqrt(radius))); + } else { // Apply the slant and a scale to space so the viewport is still fully covered by the geometry uv = vec2( - (uv.x - 0.5) * config.slantVec.x + (uv.y - 0.5) * -config.slantVec.y, - (uv.y - 0.5) * config.slantVec.x - (uv.x - 0.5) * -config.slantVec.y + (uv.x - 0.5) * config.slantVec.x + (uv.y - 0.5) * config.slantVec.y, + (uv.y - 0.5) * config.slantVec.x - (uv.x - 0.5) * config.slantVec.y ) * config.slantScale + 0.5; } uv.y = uv.y / config.glyphHeightToWidth; @@ -410,11 +396,7 @@ fn getSymbolUV(glyphCycle : f32) -> vec2 { } else { var gridCoord : vec2 = vec2(uv * config.gridSize); var gridIndex = gridCoord.y * i32(config.gridSize.x) + gridCoord.x; - if (time.frames % 2 == 0) { - glyph = cellsPing_RO.cells[gridIndex]; - } else { - glyph = cellsPong_RO.cells[gridIndex]; - } + glyph = cells_RO.cells[gridIndex]; } var brightness = glyph.r; var symbolUV = getSymbolUV(glyph.g); @@ -429,6 +411,7 @@ fn getSymbolUV(glyphCycle : f32) -> vec2 { // resolve UV to cropped position of glyph in MSDF texture var glyphUV = fract(uv * config.gridSize); + glyphUV.y = 1.0 - glyphUV.y; // WebGL -> WebGPU y-flip glyphUV = glyphUV - 0.5; glyphUV = glyphUV * clamp(1.0 - config.glyphEdgeCrop, 0.0, 1.0); glyphUV = glyphUV + 0.5;