diff --git a/README.md b/README.md index 4103e3b5..65b32b56 100644 --- a/README.md +++ b/README.md @@ -1,27 +1,68 @@ WebGL Forward+ and Clustered Deferred Shading ====================== +* Ruben Young +* Tested on: **Google Chrome 143.0.7479.0 (Official Build) canary (64-bit)** on + Windows 11 Version 24H2, AMD Ryzen 7 7800X3D, RTX 4080 SUPER + **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 4** -* (TODO) YOUR NAME HERE -* Tested on: (TODO) **Google Chrome 222.2** on - Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +## Description +This WebGPU-based application is built to demonstrate the large performance difference between forward (naive), forward+, and clustered-lights deferred rendering using the CryTek Sponza demo scene. ### Live Demo -[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +[Live Demo](http://rubenaryo.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) + +![](img/sponza_4k_2048.png) + +### Forward+ +Forward+ builds on the traditional rendering pipeline by introducing tiled (or clustered) light culling. + +Forward+ divides the scene according to view-space frusta, with each frustum maintaining a list of lights that affect it. At shading time, we now only have to calculate the contributions of the lights in our local frustum. This produces markedly better results than having to iterate over every light in the scene. + +This application further divides each frustum along the view-space Z-direction into clusters with log-depth widths. Because of this, we achieve higher resolution at points closer to the camera where detail matters more due to perspective projection. + +| ![](img/clustered.png) | +|:--:| +| Example of tiled space partitioning | + +### Deferred Rendering +We also demonstrate how we can further optimize the scene by using deferred rendering. By separating the render pass into a g-buffer and a shading pass, we can reduce the waste of shading expensive fragments that will be drawn over later by instead only shading those that will be visible. + +This is especially potent for this demo due to the high number of lights making the shading of each fragment much more expensive. + +Resulting textures from the g-buffer pass: +| ![](img/albedo_gbuffer.png) | ![](img/normal_gbuffer.png) | ![](img/position_gbuffer.png) | +|:--:|:--:|:--:| +|Albedo|Normal|World-space position + +### Examples + +| ![](img/sponza_naive_2048.gif) | ![](img/sponza_deferred_2048.gif) | ![](img/sponza_deferred_5096.gif) | +|:--:|:--:|:--:| +|Naive (2048 Lights): 10FPS|Clustered Deferred (2048 Lights): 100FPS|Clustered Deferred (4096 Lights): 240FPS| + +## Performance + +### By Light Count +Predictably, deferred shading has very good performance as the number of lights increases. Each new light adds further computation required per-fragment, so the reduction of wasted fragments gained from deferred shading is significant. + +However, we can see that deferred begins to approach forward+ at very high light counts. This is because each cluster has a hard cap on the number of lights it can contain (1024), so at very high lights we see maxed out clusters, yet lose on the performance overhead of the additional g-buffer pass. -### Demo Video/GIF +| ![](img/fpschart.png) | +|:--:| +| Note: FPS Cap is 240 due to VSync| -[![](img/video.mp4)](TODO) +### By Tile Count +Looking at the performance by tile count, we can see that there is a performance boost from further dividing the screen in XY, but we see diminishing returns at higher tile counts and hit a hard buffer size limit at 36 x 37. -### (TODO: Your README) +| ![](img/dimchart.png) | +|:--:| +| Any performance gains for Forward+ disappear after about 16x9 | -*DO NOT* leave the README to the last minute! It is a crucial part of the -project, and we will not be able to grade you without a good README. -This assignment has a considerable amount of performance analysis compared -to implementation work. Complete the implementation early to leave time! +### ### Credits diff --git a/img/albedo_gbuffer.png b/img/albedo_gbuffer.png new file mode 100644 index 00000000..c251bb93 Binary files /dev/null and b/img/albedo_gbuffer.png differ diff --git a/img/clustered.png b/img/clustered.png new file mode 100644 index 00000000..35fe20bb Binary files /dev/null and b/img/clustered.png differ diff --git a/img/dimchart.png b/img/dimchart.png new file mode 100644 index 00000000..068b6762 Binary files /dev/null and b/img/dimchart.png differ diff --git a/img/fpschart.png b/img/fpschart.png new file mode 100644 index 00000000..d067a0de Binary files /dev/null and b/img/fpschart.png differ diff --git a/img/normal_gbuffer.png b/img/normal_gbuffer.png new file mode 100644 index 00000000..50652ccd Binary files /dev/null and b/img/normal_gbuffer.png differ diff --git a/img/position_gbuffer.png b/img/position_gbuffer.png new file mode 100644 index 00000000..bdff397a Binary files /dev/null and b/img/position_gbuffer.png differ diff --git a/img/sponza_4k_2048.png b/img/sponza_4k_2048.png new file mode 100644 index 00000000..a5944bdd Binary files /dev/null and b/img/sponza_4k_2048.png differ diff --git a/img/sponza_deferred_2048.gif b/img/sponza_deferred_2048.gif new file mode 100644 index 00000000..017e9a89 Binary files /dev/null and b/img/sponza_deferred_2048.gif differ diff --git a/img/sponza_deferred_2048.mp4 b/img/sponza_deferred_2048.mp4 new file mode 100644 index 00000000..3f327d1b Binary files /dev/null and b/img/sponza_deferred_2048.mp4 differ diff --git a/img/sponza_deferred_5096.gif b/img/sponza_deferred_5096.gif new file mode 100644 index 00000000..b9883b3c Binary files /dev/null and b/img/sponza_deferred_5096.gif differ diff --git a/img/sponza_deferred_5096.mp4 b/img/sponza_deferred_5096.mp4 new file mode 100644 index 00000000..9e44d86c Binary files /dev/null and b/img/sponza_deferred_5096.mp4 differ diff --git a/img/sponza_naive_2048.gif b/img/sponza_naive_2048.gif new file mode 100644 index 00000000..73e3fc05 Binary files /dev/null and b/img/sponza_naive_2048.gif differ diff --git a/img/sponza_naive_2048.mp4 b/img/sponza_naive_2048.mp4 new file mode 100644 index 00000000..5c2dadfe Binary files /dev/null and b/img/sponza_naive_2048.mp4 differ diff --git a/src/main.ts b/src/main.ts index fa689254..294bd191 100644 --- a/src/main.ts +++ b/src/main.ts @@ -50,7 +50,7 @@ function setRenderer(mode: string) { } const renderModes = { naive: 'naive', forwardPlus: 'forward+', clusteredDeferred: 'clustered deferred' }; -let renderModeController = gui.add({ mode: renderModes.naive }, 'mode', renderModes); +let renderModeController = gui.add({ mode: renderModes.clusteredDeferred }, 'mode', renderModes); renderModeController.onChange(setRenderer); setRenderer(renderModeController.getValue()); diff --git a/src/renderers/clustered_deferred.ts b/src/renderers/clustered_deferred.ts index 00a326ca..c9053efb 100644 --- a/src/renderers/clustered_deferred.ts +++ b/src/renderers/clustered_deferred.ts @@ -2,21 +2,350 @@ import * as renderer from '../renderer'; import * as shaders from '../shaders/shaders'; import { Stage } from '../stage/stage'; -export class ClusteredDeferredRenderer extends renderer.Renderer { - // TODO-3: add layouts, pipelines, textures, etc. needed for Forward+ here +const ALBEDO_TEXTURE_FORMAT = 'rgba8unorm'; +const NORMAL_TEXTURE_FORMAT = 'rgba16float'; +const POSITION_TEXTURE_FORMAT = 'rgba16float'; +const DEPTH_TEXTURE_FORMAT = 'depth24plus'; + +export class ClusteredDeferredRenderer extends renderer.Renderer +{ + // 3: add layouts, pipelines, textures, etc. needed for Forward+ here // you may need extra uniforms such as the camera view matrix and the canvas resolution + gBufferBindGroupLayout: GPUBindGroupLayout; + gBufferBindGroup: GPUBindGroup; + + shadingBindGroupLayout: GPUBindGroupLayout; + shadingBindGroup: GPUBindGroup; + + albedoTexture: GPUTexture; + albedoTextureView: GPUTextureView; + + positionTexture: GPUTexture; + positionTextureView: GPUTextureView; + + normalTexture: GPUTexture; + normalTextureView: GPUTextureView; + + depthTexture: GPUTexture; + depthTextureView: GPUTextureView; + + sampler: GPUSampler; + + gBufferPipeline: GPURenderPipeline; + shadingPipeline: GPURenderPipeline; + constructor(stage: Stage) { super(stage); - // TODO-3: initialize layouts, pipelines, textures, etc. needed for Forward+ here - // you'll need two pipelines: one for the G-buffer pass and one for the fullscreen pass + // Textures + this.albedoTexture = renderer.device.createTexture({ + size: [renderer.canvas.width, renderer.canvas.height], + format: ALBEDO_TEXTURE_FORMAT, + usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.RENDER_ATTACHMENT + }); + this.albedoTextureView = this.albedoTexture.createView(); + + this.positionTexture = renderer.device.createTexture({ + size: [renderer.canvas.width, renderer.canvas.height], + format: POSITION_TEXTURE_FORMAT, + usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.RENDER_ATTACHMENT + }); + this.positionTextureView = this.positionTexture.createView(); + + this.normalTexture = renderer.device.createTexture({ + size: [renderer.canvas.width, renderer.canvas.height], + format: NORMAL_TEXTURE_FORMAT, + usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.RENDER_ATTACHMENT + }); + this.normalTextureView = this.normalTexture.createView(); + + this.depthTexture = renderer.device.createTexture({ + size: [renderer.canvas.width, renderer.canvas.height], + format: DEPTH_TEXTURE_FORMAT, + usage: GPUTextureUsage.RENDER_ATTACHMENT + }); + this.depthTextureView = this.depthTexture.createView(); + + // Sampler + let samplerDescriptor: GPUSamplerDescriptor = {}; + samplerDescriptor.magFilter = "linear"; + samplerDescriptor.minFilter = "linear"; + samplerDescriptor.mipmapFilter = "linear"; + samplerDescriptor.addressModeU = "repeat"; + samplerDescriptor.addressModeV = "repeat"; + + this.sampler = renderer.device.createSampler(samplerDescriptor); + + // Bind Groups + this.gBufferBindGroupLayout = renderer.device.createBindGroupLayout({ + label: "g-buffer pass bind group layout", + entries: [ + { // camUniforms + binding: 0, + visibility: GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT, + buffer: {type: "uniform"} + } + ] + }); + + this.gBufferBindGroup = renderer.device.createBindGroup({ + label: "g-buffer pass bind group", + layout: this.gBufferBindGroupLayout, + entries: [ + { + binding: 0, + resource: { buffer: this.camera.uniformsBuffer } + } + ] + }); + + this.shadingBindGroupLayout = renderer.device.createBindGroupLayout({ + label: "shading pass bind group layout", + entries: [ + { // lightSet + binding: 0, + visibility: GPUShaderStage.FRAGMENT, + buffer: { type: "read-only-storage" } + }, + { // clusterSet + binding: 1, + visibility: GPUShaderStage.FRAGMENT, + buffer: {type: "read-only-storage" } + }, + // Textures for deferred rendering + { // albedo + binding: 2, + visibility: GPUShaderStage.FRAGMENT, + texture: { + sampleType: 'float' + } + }, + { // position + binding: 3, + visibility: GPUShaderStage.FRAGMENT, + texture: { + sampleType: 'float' + } + }, + { // normal + binding: 4, + visibility: GPUShaderStage.FRAGMENT, + texture: { + sampleType: 'float' + } + }, + { // textureSampler + binding: 5, + visibility: GPUShaderStage.FRAGMENT, + sampler: {} + } + ] + }); + + this.shadingBindGroup = renderer.device.createBindGroup({ + label: "shading pass bind group", + layout: this.shadingBindGroupLayout, + entries: [ + { + binding: 0, + resource: { buffer: this.lights.lightSetStorageBuffer } + }, + { + binding: 1, + resource: { buffer: this.lights.clusterSetStorageBuffer } + }, + // Textures for Deferred rendering + { + binding: 2, + resource: this.albedoTextureView + }, + { + binding: 3, + resource: this.positionTextureView + }, + { + binding: 4, + resource: this.normalTextureView + }, + { + binding: 5, + resource: this.sampler + } + ] + }); + + + // Pipelines + this.gBufferPipeline = renderer.device.createRenderPipeline({ + layout: renderer.device.createPipelineLayout({ + label: "deferred g-buffer pipeline layout", + bindGroupLayouts: [ + this.gBufferBindGroupLayout, + renderer.modelBindGroupLayout, + renderer.materialBindGroupLayout + ] + }), + depthStencil: { + depthWriteEnabled: true, + depthCompare: "less", + format: DEPTH_TEXTURE_FORMAT + }, + vertex: { + module: renderer.device.createShaderModule({ + label: "naive vert shader", // no fplus vs + code: shaders.naiveVertSrc + }), + buffers: [ renderer.vertexBufferLayout ] + }, + fragment: { + module: renderer.device.createShaderModule({ + label: "deferred frag shader", + code: shaders.clusteredDeferredFragSrc + }), + targets: [ + { + format: ALBEDO_TEXTURE_FORMAT, + }, + { + format: POSITION_TEXTURE_FORMAT, + }, + { + format: NORMAL_TEXTURE_FORMAT, + } + ] + } + }); + + this.shadingPipeline = renderer.device.createRenderPipeline({ + layout: renderer.device.createPipelineLayout({ + label: "deferred shading pipeline layout", + bindGroupLayouts: [ + this.gBufferBindGroupLayout, + this.shadingBindGroupLayout + ] + }), + vertex: { + module: renderer.device.createShaderModule({ + label: "clustered deferred full screen vs", // no fplus vs + code: shaders.clusteredDeferredFullscreenVertSrc + }), + buffers: [ ] + }, + fragment: { + module: renderer.device.createShaderModule({ + label: "clustered deferred full screen fs", + code: shaders.clusteredDeferredFullscreenFragSrc + }), + targets: [ + { + format: renderer.canvasFormat, + } + ] + } + }); + } + + doLightClustering() + { + let encoder = renderer.device.createCommandEncoder(); + this.lights.doLightClustering(encoder); + renderer.device.queue.submit([encoder.finish()]); + } + + getGBufferRenderPassDescriptor() + { + let gbufferRenderPassDesc: GPURenderPassDescriptor = { + label: "g-buffer pass descriptor", + colorAttachments: [ + { + view: this.albedoTextureView, + clearValue: {r:0, g:0, b:0, a:1}, + loadOp: "clear", + storeOp: "store" + }, + { + view: this.positionTextureView, + clearValue: {r:0, g:0, b:0, a:0}, + loadOp: "clear", + storeOp: "store" + }, + { + view: this.normalTextureView, + clearValue: {r:0, g:0, b:0, a:0}, + loadOp: "clear", + storeOp: "store" + } + ], + depthStencilAttachment: { + view: this.depthTextureView, + depthClearValue: 1.0, + depthLoadOp: "clear", + depthStoreOp: "store" + } + }; + + return gbufferRenderPassDesc; + } + + getShadingRenderPassDescriptor() + { + let shadingRenderPassDesc: GPURenderPassDescriptor = { + label: "shading pass descriptor", + colorAttachments: [ + { + view: renderer.context.getCurrentTexture().createView(), + clearValue: [0, 0, 0, 0], + loadOp: "clear", + storeOp: "store" + } + ] + }; + return shadingRenderPassDesc + } + + doGBufferPass(pipeline: GPURenderPipeline, passDescriptor: GPURenderPassDescriptor) + { + let encoder = renderer.device.createCommandEncoder(); + const renderPass = encoder.beginRenderPass(passDescriptor); + renderPass.setPipeline(pipeline); + + renderPass.setBindGroup(shaders.constants.bindGroup_scene, this.gBufferBindGroup); + + this.scene.iterate(node => { + renderPass.setBindGroup(shaders.constants.bindGroup_model, node.modelBindGroup); + }, material => { + renderPass.setBindGroup(shaders.constants.bindGroup_material, material.materialBindGroup); + }, primitive => { + renderPass.setVertexBuffer(0, primitive.vertexBuffer); + renderPass.setIndexBuffer(primitive.indexBuffer, 'uint32'); + renderPass.drawIndexed(primitive.numIndices); + }); + + renderPass.end(); + + renderer.device.queue.submit([encoder.finish()]); + } + + doFullscreenPass(pipeline: GPURenderPipeline, passDescriptor: GPURenderPassDescriptor) + { + let encoder = renderer.device.createCommandEncoder(); + const renderPass = encoder.beginRenderPass(passDescriptor); + + // Fullscreen pass is simple, only the basic bind groups and drawing only 6 verts for the screen quad. + renderPass.setPipeline(pipeline); + renderPass.setBindGroup(0, this.gBufferBindGroup); + renderPass.setBindGroup(1, this.shadingBindGroup); + renderPass.draw(6); + renderPass.end(); + + renderer.device.queue.submit([encoder.finish()]); } - override draw() { - // TODO-3: run the Forward+ rendering pass: - // - run the clustering compute shader - // - run the G-buffer pass, outputting position, albedo, and normals - // - run the fullscreen pass, which reads from the G-buffer and performs lighting calculations + override draw() + { + this.doLightClustering(); + this.doGBufferPass(this.gBufferPipeline, this.getGBufferRenderPassDescriptor()); + this.doFullscreenPass(this.shadingPipeline, this.getShadingRenderPassDescriptor()); } -} +} \ No newline at end of file diff --git a/src/renderers/forward_plus.ts b/src/renderers/forward_plus.ts index 471796fd..8c946c14 100644 --- a/src/renderers/forward_plus.ts +++ b/src/renderers/forward_plus.ts @@ -5,16 +5,143 @@ import { Stage } from '../stage/stage'; export class ForwardPlusRenderer extends renderer.Renderer { // TODO-2: add layouts, pipelines, textures, etc. needed for Forward+ here // you may need extra uniforms such as the camera view matrix and the canvas resolution + + sceneUniformsBindGroupLayout: GPUBindGroupLayout; + sceneUniformsBindGroup: GPUBindGroup; + + depthTexture: GPUTexture; + depthTextureView: GPUTextureView; + + pipeline: GPURenderPipeline; constructor(stage: Stage) { super(stage); - // TODO-2: initialize layouts, pipelines, textures, etc. needed for Forward+ here + this.sceneUniformsBindGroupLayout = renderer.device.createBindGroupLayout({ + label: "scene uniforms bind group layout", + entries: [ + { // camUniforms + binding: 0, + visibility: GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT, + buffer: {type: "uniform"} + }, + { // lightSet + binding: 1, + visibility: GPUShaderStage.FRAGMENT, + buffer: { type: "read-only-storage" } + }, + { // clusterSet + binding: 2, + visibility: GPUShaderStage.FRAGMENT, + buffer: {type: "read-only-storage" } + } + ] + }); + + this.sceneUniformsBindGroup = renderer.device.createBindGroup({ + label: "scene uniforms bind group", + layout: this.sceneUniformsBindGroupLayout, + entries: [ + // 1.2: add an entry for camera uniforms at binding 0 + // you can access the camera using `this.camera` + // if you run into TypeScript errors, you're probably trying to upload the host buffer instead + { + binding: 0, + resource: { buffer: this.camera.uniformsBuffer } + }, + { + binding: 1, + resource: { buffer: this.lights.lightSetStorageBuffer } + }, + { + binding: 2, + resource: { buffer: this.lights.clusterSetStorageBuffer } + } + ] + }); + + this.depthTexture = renderer.device.createTexture({ + size: [renderer.canvas.width, renderer.canvas.height], + format: "depth24plus", + usage: GPUTextureUsage.RENDER_ATTACHMENT + }); + this.depthTextureView = this.depthTexture.createView(); + + this.pipeline = renderer.device.createRenderPipeline({ + layout: renderer.device.createPipelineLayout({ + label: "f-plus pipeline layout", + bindGroupLayouts: [ + this.sceneUniformsBindGroupLayout, + renderer.modelBindGroupLayout, + renderer.materialBindGroupLayout + ] + }), + depthStencil: { + depthWriteEnabled: true, + depthCompare: "less", + format: "depth24plus" + }, + vertex: { + module: renderer.device.createShaderModule({ + label: "naive vert shader", // no fplus vs + code: shaders.naiveVertSrc + }), + buffers: [ renderer.vertexBufferLayout ] + }, + fragment: { + module: renderer.device.createShaderModule({ + label: "f-plus frag shader", + code: shaders.forwardPlusFragSrc, + }), + targets: [ + { + format: renderer.canvasFormat, + } + ] + } + }); } override draw() { - // TODO-2: run the Forward+ rendering pass: - // - run the clustering compute shader - // - run the main rendering pass, using the computed clusters for efficient lighting + const encoder = renderer.device.createCommandEncoder(); + const canvasTextureView = renderer.context.getCurrentTexture().createView(); + + this.lights.doLightClustering(encoder); + + const renderPass = encoder.beginRenderPass({ + label: "f-plus render pass", + colorAttachments: [ + { + view: canvasTextureView, + clearValue: [0, 0, 0, 0], + loadOp: "clear", + storeOp: "store" + } + ], + depthStencilAttachment: { + view: this.depthTextureView, + depthClearValue: 1.0, + depthLoadOp: "clear", + depthStoreOp: "store" + } + }); + renderPass.setPipeline(this.pipeline); + + // 1.2: bind `this.sceneUniformsBindGroup` to index `shaders.constants.bindGroup_scene` + renderPass.setBindGroup(shaders.constants.bindGroup_scene, this.sceneUniformsBindGroup); + + this.scene.iterate(node => { + renderPass.setBindGroup(shaders.constants.bindGroup_model, node.modelBindGroup); + }, material => { + renderPass.setBindGroup(shaders.constants.bindGroup_material, material.materialBindGroup); + }, primitive => { + renderPass.setVertexBuffer(0, primitive.vertexBuffer); + renderPass.setIndexBuffer(primitive.indexBuffer, 'uint32'); + renderPass.drawIndexed(primitive.numIndices); + }); + + renderPass.end(); + + renderer.device.queue.submit([encoder.finish()]); } } diff --git a/src/renderers/naive.ts b/src/renderers/naive.ts index 0bf82417..5cb1ef74 100644 --- a/src/renderers/naive.ts +++ b/src/renderers/naive.ts @@ -17,7 +17,12 @@ export class NaiveRenderer extends renderer.Renderer { this.sceneUniformsBindGroupLayout = renderer.device.createBindGroupLayout({ label: "scene uniforms bind group layout", entries: [ - // TODO-1.2: add an entry for camera uniforms at binding 0, visible to only the vertex shader, and of type "uniform" + // 1.2: add an entry for camera uniforms at binding 0, visible to only the vertex shader, and of type "uniform" + { + binding: 0, + visibility: GPUShaderStage.VERTEX, + buffer: {type: "uniform"} + }, { // lightSet binding: 1, visibility: GPUShaderStage.FRAGMENT, @@ -30,9 +35,13 @@ export class NaiveRenderer extends renderer.Renderer { label: "scene uniforms bind group", layout: this.sceneUniformsBindGroupLayout, entries: [ - // TODO-1.2: add an entry for camera uniforms at binding 0 + // 1.2: add an entry for camera uniforms at binding 0 // you can access the camera using `this.camera` // if you run into TypeScript errors, you're probably trying to upload the host buffer instead + { + binding: 0, + resource: { buffer: this.camera.uniformsBuffer } + }, { binding: 1, resource: { buffer: this.lights.lightSetStorageBuffer } @@ -105,7 +114,8 @@ export class NaiveRenderer extends renderer.Renderer { }); renderPass.setPipeline(this.pipeline); - // TODO-1.2: bind `this.sceneUniformsBindGroup` to index `shaders.constants.bindGroup_scene` + // 1.2: bind `this.sceneUniformsBindGroup` to index `shaders.constants.bindGroup_scene` + renderPass.setBindGroup(shaders.constants.bindGroup_scene, this.sceneUniformsBindGroup); this.scene.iterate(node => { renderPass.setBindGroup(shaders.constants.bindGroup_model, node.modelBindGroup); diff --git a/src/shaders/clustered_deferred.fs.wgsl b/src/shaders/clustered_deferred.fs.wgsl index 4e86f573..6e25c4e4 100644 --- a/src/shaders/clustered_deferred.fs.wgsl +++ b/src/shaders/clustered_deferred.fs.wgsl @@ -1,3 +1,37 @@ // TODO-3: implement the Clustered Deferred G-buffer fragment shader // This shader should only store G-buffer information and should not do any shading. +@group(${bindGroup_material}) @binding(0) var diffuseTex: texture_2d; +@group(${bindGroup_material}) @binding(1) var diffuseTexSampler: sampler; + +struct FragmentInput +{ + @location(0) pos: vec3f, + @location(1) nor: vec3f, + @location(2) uv: vec2f +} + +struct GBufferOut +{ + @location(0) albedo: vec4, + @location(1) position: vec4, + @location(2) normal: vec4, +} + +@fragment +fn main(in: FragmentInput) -> GBufferOut +{ + let diffuseColor = textureSample(diffuseTex, diffuseTexSampler, in.uv); + if (diffuseColor.a < 0.5) { + discard; + } + + let normalRGB = (in.nor * 0.5) + vec3(0.5); + + return GBufferOut + ( + diffuseColor, + vec4f(in.pos, 1.0), + vec4f(normalRGB.rgb, 0.0) + ); +} diff --git a/src/shaders/clustered_deferred_fullscreen.fs.wgsl b/src/shaders/clustered_deferred_fullscreen.fs.wgsl index 68235c41..efb169a0 100644 --- a/src/shaders/clustered_deferred_fullscreen.fs.wgsl +++ b/src/shaders/clustered_deferred_fullscreen.fs.wgsl @@ -1,3 +1,56 @@ -// TODO-3: implement the Clustered Deferred fullscreen fragment shader +@group(0) @binding(0) var camUniforms: CameraUniforms; -// Similar to the Forward+ fragment shader, but with vertex information coming from the G-buffer instead. +@group(1) @binding(0) var lightSet: LightSet; +@group(1) @binding(1) var clusterSet: ClusterSet; +@group(1) @binding(2) var albedoTexture: texture_2d; +@group(1) @binding(3) var positionTexture: texture_2d; +@group(1) @binding(4) var normalTexture: texture_2d; +@group(1) @binding(5) var textureSampler: sampler; + +struct FragmentInput +{ + @builtin(position) pos : vec4, +} + +@fragment +fn main(in: FragmentInput) -> @location(0) vec4f +{ + let uv = in.pos.xy / camUniforms.resolution.xy; + let albedo = textureSample(albedoTexture, textureSampler, uv); + let position = textureSample(positionTexture, textureSampler, uv); + var normal = textureSample(normalTexture, textureSampler, uv); + + normal *= 2.0; + normal -= vec4f(1.0); + normal.w = 0.0; + + const numClustersX = ${clusterCountX}; + const numClustersY = ${clusterCountY}; + const numClustersZ = ${clusterCountZ}; + + let clusterIdxX = u32(in.pos.x / f32(camUniforms.resolution.x) * f32(numClustersX)); + let clusterIdxY = u32(in.pos.y / f32(camUniforms.resolution.y) * f32(numClustersY)); + + // Calculate Z using log depth formula + let viewSpacePos = camUniforms.view * vec4f(position.xyz, 1.0); + let viewDepth = viewSpacePos.z; + let logDepth = log(-viewDepth / camUniforms.near) / log(camUniforms.far / camUniforms.near); + + let clusterIdxZ = clamp(u32(logDepth * numClustersZ), 0u, u32(numClustersZ - 1u)); + + // Only check lights that are in this cluster + let clusterIndex = u32(clusterIdxX + (clusterIdxY * numClustersX) + clusterIdxZ * (numClustersX * numClustersY)); + let numLights = u32(clusterSet.clusters[clusterIndex].numLights); + + const AMBIENT_LIGHT = 0.00; + var totalLightContrib = vec3f(AMBIENT_LIGHT); + for (var lightIdx = 0u; lightIdx < numLights; lightIdx++) { + + let mainLightIndex = clusterSet.clusters[clusterIndex].lights[lightIdx]; + let light = lightSet.lights[mainLightIndex]; + totalLightContrib += calculateLightContrib(light, position.xyz, normalize(normal.xyz)); + } + + var finalColor = albedo.rgb * totalLightContrib; + return vec4f(finalColor, 1.0); +} \ No newline at end of file diff --git a/src/shaders/clustered_deferred_fullscreen.vs.wgsl b/src/shaders/clustered_deferred_fullscreen.vs.wgsl index 1e43a884..128fe872 100644 --- a/src/shaders/clustered_deferred_fullscreen.vs.wgsl +++ b/src/shaders/clustered_deferred_fullscreen.vs.wgsl @@ -1,3 +1,25 @@ -// TODO-3: implement the Clustered Deferred fullscreen vertex shader +// 3: implement the Clustered Deferred fullscreen vertex shader // This shader should be very simple as it does not need all of the information passed by the the naive vertex shader. +struct VertexOutput +{ + @builtin(position) pos : vec4, +} + +@vertex +fn main(@builtin(vertex_index) vertIdx: u32) -> VertexOutput { + + // full screen quad + var vbo = array, 6> + ( + vec2(-1.0, -1.0), + vec2( 1.0, -1.0), + vec2(-1.0, 1.0), + + vec2(-1.0, 1.0), + vec2( 1.0, -1.0), + vec2( 1.0, 1.0), + ); + + return VertexOutput(vec4(vbo[vertIdx].xy, 0.0, 1.0)); +} \ No newline at end of file diff --git a/src/shaders/clustering.cs.wgsl b/src/shaders/clustering.cs.wgsl index 575d6e5a..91dc3637 100644 --- a/src/shaders/clustering.cs.wgsl +++ b/src/shaders/clustering.cs.wgsl @@ -1,4 +1,6 @@ -// TODO-2: implement the light clustering compute shader +@group(${bindGroup_scene}) @binding(0) var camUniforms: CameraUniforms; +@group(${bindGroup_scene}) @binding(1) var lightSet: LightSet; +@group(${bindGroup_scene}) @binding(2) var clusterSet: ClusterSet; // ------------------------------------ // Calculating cluster bounds: @@ -21,3 +23,98 @@ // - Stop adding lights if the maximum number of lights is reached. // - Store the number of lights assigned to this cluster. + +fn clipToView(clip : vec4f) -> vec4f +{ + let invProj = camUniforms.invProj; + + var view = invProj * clip; + view = view / view.w; + + return view; +} + +fn pixelToView(pixel: vec4f) -> vec4f +{ + let uv = (pixel.xy / camUniforms.resolution); + let clip = vec4f(vec2f(uv.x, 1.0 - uv.y), pixel.z, pixel.w) * vec4f(2.0) - vec4f(1.0); + return clipToView(clip); +} + +fn lightIntersection(clusterMin: vec3f, clusterMax: vec3f, lightPosView: vec3f) -> bool +{ + // If the bounding region for the cluster has any contribution from the light, return true + let boundaryPoint = clamp(lightPosView, clusterMin, clusterMax); + + let lightToBoundary = lightPosView - boundaryPoint; + let sqDist = f32(dot(lightToBoundary, lightToBoundary)); + + // If the distance from the light to the closest point is less than the radius, this light affects the AABB. + return sqDist < f32(LIGHT_RADIUS * LIGHT_RADIUS); +} + +@compute +@workgroup_size(${clusterWorkgroupDimX}, ${clusterWorkgroupDimY}, ${clusterWorkgroupDimZ}) +fn main(@builtin(global_invocation_id) globalIdx: vec3u) +{ + const numClustersX = ${clusterCountX}; + const numClustersY = ${clusterCountY}; + const numClustersZ = ${clusterCountZ}; + + if (globalIdx.x >= numClustersX || globalIdx.y >= numClustersY || globalIdx.z >= numClustersZ) + { + return; + } + + // 3D -> 1D index + let clusterIndex = globalIdx.x + (globalIdx.y * numClustersX) + globalIdx.z * (numClustersX * numClustersY); + + let resolution = camUniforms.resolution; + let tileSizeX = resolution.x / numClustersX; + let tileSizeY = resolution.y / numClustersY; + + // In Pixel Space + let clusterMinX = f32(globalIdx.x) * tileSizeX; + let clusterMaxX = f32(globalIdx.x+1) * tileSizeX; + + let clusterMinY = f32(globalIdx.y) * tileSizeY; + let clusterMaxY = f32(globalIdx.y+1) * tileSizeY; + + let clusterMinXMinY_pixel = vec4f(clusterMinX, clusterMinY, -1.0, 1.0); + let clusterMaxXMaxY_pixel = vec4f(clusterMaxX, clusterMaxY, -1.0, 1.0); + + // Convert to View space + let clusterMin_view = pixelToView(clusterMinXMinY_pixel).xyz; + let clusterMax_view = pixelToView(clusterMaxXMaxY_pixel).xyz; + + // Compute z bounds by log depth + let near = camUniforms.near; + let far = camUniforms.far; + let clusterNear = -near * pow(far/near, f32(globalIdx.z) / numClustersZ); + let clusterFar = -near * pow(far/near, f32(globalIdx.z+1) / numClustersZ); + + let minPointNear = clusterMin_view * (clusterNear / clusterMin_view.z); + let minPointFar = clusterMin_view * (clusterFar / clusterMin_view.z); + let maxPointNear = clusterMax_view * (clusterNear / clusterMax_view.z); + let maxPointFar = clusterMax_view * (clusterFar / clusterMax_view.z); + + let clusterMin = min(min(minPointNear, minPointFar),min(maxPointNear, maxPointFar)); + let clusterMax = max(max(minPointNear, minPointFar),max(maxPointNear, maxPointFar)); + + // Iterate over all the lights and determine if they would affect this cluster + clusterSet.clusters[clusterIndex].numLights = 0u; + var numLightsInCluster = 0u; + for (var lightIdx = 0u; lightIdx < lightSet.numLights && numLightsInCluster < ${maxLightsPerCluster}; lightIdx++) + { + let light = lightSet.lights[lightIdx]; + let lightPosView = camUniforms.view * vec4(light.pos, 1.0); + + if (lightIntersection(clusterMin, clusterMax, lightPosView.xyz)) + { + // This light affects the cluster + clusterSet.clusters[clusterIndex].lights[numLightsInCluster] = lightIdx; + numLightsInCluster++; + } + } + clusterSet.clusters[clusterIndex].numLights = numLightsInCluster; +} diff --git a/src/shaders/common.wgsl b/src/shaders/common.wgsl index 738e9c4e..f96c4964 100644 --- a/src/shaders/common.wgsl +++ b/src/shaders/common.wgsl @@ -10,15 +10,34 @@ struct LightSet { lights: array } -// TODO-2: you may want to create a ClusterSet struct similar to LightSet +// 2: you may want to create a ClusterSet struct similar to LightSet +struct Cluster +{ + numLights: u32, + lights: array +} + +struct ClusterSet +{ + count: u32, + clusters: array +} struct CameraUniforms { - // TODO-1.3: add an entry for the view proj mat (of type mat4x4f) + // 1.3: add an entry for the view proj mat (of type mat4x4f) + view: mat4x4f, + viewProj: mat4x4f, + invProj: mat4x4f, + resolution: vec2f, + near: f32, + far: f32 } +const LIGHT_RADIUS: f32 = 2.0; + // CHECKITOUT: this special attenuation function ensures lights don't affect geometry outside the maximum light radius fn rangeAttenuation(distance: f32) -> f32 { - return clamp(1.f - pow(distance / ${lightRadius}, 4.f), 0.f, 1.f) / (distance * distance); + return clamp(1.f - pow(distance / LIGHT_RADIUS, 4.f), 0.f, 1.f) / (distance * distance); } fn calculateLightContrib(light: Light, posWorld: vec3f, nor: vec3f) -> vec3f { @@ -27,4 +46,4 @@ fn calculateLightContrib(light: Light, posWorld: vec3f, nor: vec3f) -> vec3f { let lambert = max(dot(nor, normalize(vecToLight)), 0.f); return light.color * lambert * rangeAttenuation(distToLight); -} +} \ No newline at end of file diff --git a/src/shaders/forward_plus.fs.wgsl b/src/shaders/forward_plus.fs.wgsl index 0500e3df..0aeabc10 100644 --- a/src/shaders/forward_plus.fs.wgsl +++ b/src/shaders/forward_plus.fs.wgsl @@ -1,4 +1,4 @@ -// TODO-2: implement the Forward+ fragment shader +// 2: implement the Forward+ fragment shader // See naive.fs.wgsl for basic fragment shader setup; this shader should use light clusters instead of looping over all lights @@ -14,3 +14,57 @@ // Add the calculated contribution to the total light accumulation. // Multiply the fragment’s diffuse color by the accumulated light contribution. // Return the final color, ensuring that the alpha component is set appropriately (typically to 1). + +@group(${bindGroup_scene}) @binding(0) var camUniforms: CameraUniforms; +@group(${bindGroup_scene}) @binding(1) var lightSet: LightSet; +@group(${bindGroup_scene}) @binding(2) var clusterSet: ClusterSet; + +@group(${bindGroup_material}) @binding(0) var diffuseTex: texture_2d; +@group(${bindGroup_material}) @binding(1) var diffuseTexSampler: sampler; + +struct FragmentInput +{ + @location(0) pos: vec3f, + @location(1) nor: vec3f, + @location(2) uv: vec2f +} + +@fragment +fn main(in: FragmentInput, @builtin(position) fragCoord: vec4f) -> @location(0) vec4f +{ + let diffuseColor = textureSample(diffuseTex, diffuseTexSampler, in.uv); + if (diffuseColor.a < 0.5) + { + discard; + } + + const numClustersX = ${clusterCountX}; + const numClustersY = ${clusterCountY}; + const numClustersZ = ${clusterCountZ}; + + let clusterIdxX = u32(f32(fragCoord.x) / f32(camUniforms.resolution.x) * f32(numClustersX)); + let clusterIdxY = u32(f32(fragCoord.y) / f32(camUniforms.resolution.y) * f32(numClustersY)); + + // Calculate Z using log depth formula + let viewSpacePos = camUniforms.view * vec4f(in.pos, 1.0); + let viewDepth = viewSpacePos.z; + let logDepth = log(-viewDepth / camUniforms.near) / log(camUniforms.far / camUniforms.near); + + let clusterIdxZ = clamp(u32(logDepth * numClustersZ), 0u, u32(numClustersZ - 1u)); + + // Only check lights that are in this cluster + let clusterIndex = u32(clusterIdxX + (clusterIdxY * numClustersX) + clusterIdxZ * (numClustersX * numClustersY)); + let numLights = u32(clusterSet.clusters[clusterIndex].numLights); + + const AMBIENT_LIGHT = 0.00; + var totalLightContrib = vec3f(AMBIENT_LIGHT); + for (var lightIdx = 0u; lightIdx < numLights; lightIdx++) { + + let mainLightIndex = clusterSet.clusters[clusterIndex].lights[lightIdx]; + let light = lightSet.lights[mainLightIndex]; + totalLightContrib += calculateLightContrib(light, in.pos, normalize(in.nor)); + } + + var finalColor = diffuseColor.rgb * totalLightContrib; + return vec4f(finalColor, 1.0); +} diff --git a/src/shaders/naive.vs.wgsl b/src/shaders/naive.vs.wgsl index 5a7ddd4b..aacd9d7c 100644 --- a/src/shaders/naive.vs.wgsl +++ b/src/shaders/naive.vs.wgsl @@ -1,7 +1,8 @@ // CHECKITOUT: you can use this vertex shader for all of the renderers -// TODO-1.3: add a uniform variable here for camera uniforms (of type CameraUniforms) +// 1.3: add a uniform variable here for camera uniforms (of type CameraUniforms) // make sure to use ${bindGroup_scene} for the group +@group(${bindGroup_scene}) @binding(0) var camUniforms: CameraUniforms; @group(${bindGroup_model}) @binding(0) var modelMat: mat4x4f; @@ -26,7 +27,7 @@ fn main(in: VertexInput) -> VertexOutput let modelPos = modelMat * vec4(in.pos, 1); var out: VertexOutput; - out.fragPos = ??? * modelPos; // TODO-1.3: replace ??? with the view proj mat from your CameraUniforms uniform variable + out.fragPos = camUniforms.viewProj * modelPos; // 1.3: replace ??? with the view proj mat from your CameraUniforms uniform variable out.pos = modelPos.xyz / modelPos.w; out.nor = in.nor; out.uv = in.uv; diff --git a/src/shaders/shaders.ts b/src/shaders/shaders.ts index 584c008f..a6ec9e0e 100644 --- a/src/shaders/shaders.ts +++ b/src/shaders/shaders.ts @@ -1,7 +1,4 @@ // CHECKITOUT: this file loads all the shaders and preprocesses them with some common code - -import { Camera } from '../stage/camera'; - import commonRaw from './common.wgsl?raw'; import naiveVertRaw from './naive.vs.wgsl?raw'; @@ -27,10 +24,22 @@ export const constants = { bindGroup_scene: 0, bindGroup_model: 1, bindGroup_material: 2, + bindGroup_shading: 3, moveLightsWorkgroupSize: 128, - - lightRadius: 2 + + // Cluster counts in each dimension + clusterCountX: 16, + clusterCountY: 9, + clusterCountZ: 24, + + // workgroup size for cluster rendering + // 8 * 8 * 4 = 256 + clusterWorkgroupDimX: 8, + clusterWorkgroupDimY: 8, + clusterWorkgroupDimZ: 4, + + maxLightsPerCluster: 1024 }; // ================================= diff --git a/src/stage/camera.ts b/src/stage/camera.ts index 7d2a4a1e..98843a65 100644 --- a/src/stage/camera.ts +++ b/src/stage/camera.ts @@ -2,19 +2,61 @@ import { Mat4, mat4, Vec3, vec3 } from "wgpu-matrix"; import { toRadians } from "../math_util"; import { device, canvas, fovYDegrees, aspectRatio } from "../renderer"; -class CameraUniforms { - readonly buffer = new ArrayBuffer(16 * 4); - private readonly floatView = new Float32Array(this.buffer); +const CameraUniformsValues = new ArrayBuffer(208); +class CameraUniformsViews { + private readonly view = new Float32Array(CameraUniformsValues, 0, 16); + private readonly viewProj = new Float32Array(CameraUniformsValues, 64, 16); + private readonly invProj = new Float32Array(CameraUniformsValues, 128, 16); + private readonly resolution = new Float32Array(CameraUniformsValues, 192, 2); + private readonly near = new Float32Array(CameraUniformsValues, 200, 1); + private readonly far = new Float32Array(CameraUniformsValues, 204, 1); + + setView(mat: Float32Array) + { + for (let i = 0; i != 16; ++i) + { + this.view[i] = mat[i]; + } + } + + setViewProj(mat: Float32Array) + { + // 1.1: set the first 16 elements of `this.floatView` to the input `mat` + for (let i = 0; i != 16; ++i) + { + this.viewProj[i] = mat[i]; + } + } + + setInvProj(mat: Float32Array) + { + for (let i = 0; i != 16; ++i) + { + this.invProj[i] = mat[i]; + } + } + + setResolution(w:number, h:number) + { + this.resolution[0] = w; + this.resolution[1] = h; + } - set viewProjMat(mat: Float32Array) { - // TODO-1.1: set the first 16 elements of `this.floatView` to the input `mat` + setNear(n:number) + { + this.near[0] = n; } - // TODO-2: add extra functions to set values needed for light clustering here + setFar(f:number) + { + this.far[0] = f; + } + + // 2: add extra functions to set values needed for light clustering here } export class Camera { - uniforms: CameraUniforms = new CameraUniforms(); + uniforms: CameraUniformsViews = new CameraUniformsViews(); uniformsBuffer: GPUBuffer; projMat: Mat4 = mat4.create(); @@ -33,14 +75,22 @@ export class Camera { keys: { [key: string]: boolean } = {}; constructor () { - // TODO-1.1: set `this.uniformsBuffer` to a new buffer of size `this.uniforms.buffer.byteLength` + // 1.1: set `this.uniformsBuffer` to a new buffer of size `this.uniforms.buffer.byteLength` // ensure the usage is set to `GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST` since we will be copying to this buffer // check `lights.ts` for examples of using `device.createBuffer()` // // note that you can add more variables (e.g. inverse proj matrix) to this buffer in later parts of the assignment - + this.uniformsBuffer = device.createBuffer({ + label: "uniforms", + size: CameraUniformsValues.byteLength, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST + }); this.projMat = mat4.perspective(toRadians(fovYDegrees), aspectRatio, Camera.nearPlane, Camera.farPlane); + this.uniforms.setResolution(canvas.width, canvas.height); + this.uniforms.setNear(Camera.nearPlane); + this.uniforms.setFar(Camera.farPlane); + this.rotateCamera(0, 0); // set initial camera vectors window.addEventListener('keydown', (event) => this.onKeyEvent(event, true)); @@ -128,11 +178,15 @@ export class Camera { const lookPos = vec3.add(this.cameraPos, vec3.scale(this.cameraFront, 1)); const viewMat = mat4.lookAt(this.cameraPos, lookPos, [0, 1, 0]); const viewProjMat = mat4.mul(this.projMat, viewMat); - // TODO-1.1: set `this.uniforms.viewProjMat` to the newly calculated view proj mat + // 1.1: set `this.uniforms.viewProjMat` to the newly calculated view proj mat + this.uniforms.setView(viewMat); + this.uniforms.setViewProj(viewProjMat); + this.uniforms.setInvProj(mat4.inverse(this.projMat)); - // TODO-2: write to extra buffers needed for light clustering here + // 2: write to extra buffers needed for light clustering here - // TODO-1.1: upload `this.uniforms.buffer` (host side) to `this.uniformsBuffer` (device side) + // 1.1: upload `this.uniforms.buffer` (host side) to `this.uniformsBuffer` (device side) // check `lights.ts` for examples of using `device.queue.writeBuffer()` + device.queue.writeBuffer(this.uniformsBuffer, 0, CameraUniformsValues) } } diff --git a/src/stage/lights.ts b/src/stage/lights.ts index a6eed919..3f377580 100644 --- a/src/stage/lights.ts +++ b/src/stage/lights.ts @@ -11,10 +11,8 @@ function hueToRgb(h: number) { } export class Lights { - private camera: Camera; - - numLights = 500; - static readonly maxNumLights = 5000; + numLights = 1024; + static readonly maxNumLights = 8192; static readonly numFloatsPerLight = 8; // vec3f is aligned at 16 byte boundaries static readonly lightIntensity = 0.1; @@ -22,22 +20,26 @@ export class Lights { lightsArray = new Float32Array(Lights.maxNumLights * Lights.numFloatsPerLight); lightSetStorageBuffer: GPUBuffer; + clusterSetArray: Float32Array; + clusterSetStorageBuffer: GPUBuffer; + timeUniformBuffer: GPUBuffer; moveLightsComputeBindGroupLayout: GPUBindGroupLayout; moveLightsComputeBindGroup: GPUBindGroup; moveLightsComputePipeline: GPUComputePipeline; - // TODO-2: add layouts, pipelines, textures, etc. needed for light clustering here + // 2: add layouts, pipelines, textures, etc. needed for light clustering here + clusteringComputeBindGroupLayout: GPUBindGroupLayout; + clusteringComputeBindGroup: GPUBindGroup; + clusteringComputePipeline: GPUComputePipeline; constructor(camera: Camera) { - this.camera = camera; - this.lightSetStorageBuffer = device.createBuffer({ label: "lights", size: 16 + this.lightsArray.byteLength, // 16 for numLights + padding usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST - }); + }); this.populateLightsBuffer(); this.updateLightSetUniformNumLights(); @@ -93,7 +95,70 @@ export class Lights { } }); - // TODO-2: initialize layouts, pipelines, textures, etc. needed for light clustering here + // 2: initialize layouts, pipelines, textures, etc. needed for light clustering here + const MAX_CLUSTERS = shaders.constants.clusterCountX * shaders.constants.clusterCountY * shaders.constants.clusterCountZ; + const CLUSTER_FLOAT_COUNT = shaders.constants.maxLightsPerCluster + 4; + this.clusterSetArray = new Float32Array(MAX_CLUSTERS * CLUSTER_FLOAT_COUNT); + this.clusterSetStorageBuffer = device.createBuffer({ + label: "clusterSet", + size: 16 + this.clusterSetArray.byteLength, // 16 for cluster count + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + }); + + this.clusteringComputeBindGroupLayout = device.createBindGroupLayout({ + label: "cluster lights compute bind group layout", + entries: [ + { // camera uniforms + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: {type: "uniform"} + }, + { // light set + binding: 1, + visibility: GPUShaderStage.COMPUTE, + buffer: {type: "read-only-storage"} + }, + { // cluster set + binding: 2, + visibility: GPUShaderStage.COMPUTE, + buffer: {type: "storage"} + } + ] + }); + + this.clusteringComputeBindGroup = device.createBindGroup({ + label: "cluster lights compute bind group", + layout: this.clusteringComputeBindGroupLayout, + entries: [ + { + binding: 0, + resource: { buffer: camera.uniformsBuffer } + }, + { + binding: 1, + resource: { buffer: this.lightSetStorageBuffer } + }, + { + binding: 2, + resource: { buffer: this.clusterSetStorageBuffer} + } + ] + }); + + this.clusteringComputePipeline = device.createComputePipeline({ + label: "cluster lights compute pipeline", + layout: device.createPipelineLayout({ + label: "cluster lights compute pipeline layout", + bindGroupLayouts: [ this.clusteringComputeBindGroupLayout] + }), + compute: { + module: device.createShaderModule({ + label: "cluster lights compute shader", + code: shaders.clusteringComputeSrc + }), + entryPoint: "main" + } + }); } private populateLightsBuffer() { @@ -111,8 +176,18 @@ export class Lights { } doLightClustering(encoder: GPUCommandEncoder) { - // TODO-2: run the light clustering compute pass(es) here + // 2: run the light clustering compute pass(es) here // implementing clustering here allows for reusing the code in both Forward+ and Clustered Deferred + const computePass = encoder.beginComputePass(); + computePass.setPipeline(this.clusteringComputePipeline); + computePass.setBindGroup(0, this.clusteringComputeBindGroup); + + const numWorkgroupsX = Math.ceil(shaders.constants.clusterCountX / shaders.constants.clusterWorkgroupDimX); + const numWorkgroupsY = Math.ceil(shaders.constants.clusterCountY / shaders.constants.clusterWorkgroupDimY); + const numWorkgroupsZ = Math.ceil(shaders.constants.clusterCountZ / shaders.constants.clusterWorkgroupDimZ); + + computePass.dispatchWorkgroups(numWorkgroupsX, numWorkgroupsY, numWorkgroupsZ); + computePass.end(); } // CHECKITOUT: this is where the light movement compute shader is dispatched from the host