diff --git a/IMAGES/CLUSTERS_128_128_128.png b/IMAGES/CLUSTERS_128_128_128.png new file mode 100644 index 00000000..5d9b6c07 Binary files /dev/null and b/IMAGES/CLUSTERS_128_128_128.png differ diff --git a/IMAGES/CLUSTERS_128_128_256.png b/IMAGES/CLUSTERS_128_128_256.png new file mode 100644 index 00000000..dfef3e07 Binary files /dev/null and b/IMAGES/CLUSTERS_128_128_256.png differ diff --git a/IMAGES/CLUSTERS_64_64_128.png b/IMAGES/CLUSTERS_64_64_128.png new file mode 100644 index 00000000..356c41b5 Binary files /dev/null and b/IMAGES/CLUSTERS_64_64_128.png differ diff --git a/IMAGES/Naive1k_V1.png b/IMAGES/Naive1k_V1.png new file mode 100644 index 00000000..84f9962d Binary files /dev/null and b/IMAGES/Naive1k_V1.png differ diff --git a/README.md b/README.md index 4103e3b5..054e2453 100644 --- a/README.md +++ b/README.md @@ -3,25 +3,39 @@ WebGL Forward+ and Clustered Deferred Shading **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) +* Lewis Ghrist +* Tested on: **Google Chrome 141.0.7390.108** on + Windows 11, Intel Core i7-13700H, 32.0 GB RAM ### Live Demo -[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +[LIVE DEMO:](https://siwel-cg.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred/) +--- -### Demo Video/GIF +![Naive](IMAGES/Naive1k_V1.png) +--- +## A Quick Note +If you take a look at the live demo, you will notice there are some things missing. I wasn't able to implement all the features, and some of the ones I was able to implement leave lots of room for improvement. I am going to keep working on this and have a full analysis once everything has been implemented, but for now, you are stuck with jank. However, let just look at some of the things I did do (as of 10/18/25), since in my implementation, they are more visually interesting than a plane black screen. -[![](img/video.mp4)](TODO) +# Overview +First off, what was I trying to do. As scenes become more and more complex, we need to develop new methods for handaling the increased amounts of calculations needed to actually render these complex scene. One such complexity is the number of lights in a scene. When rasterizing, for some fragment and the object it hits, we need to calculate what that pixel is going to look like, which usually involves some sort of lighting calculation. The naive approach is to simply test each light for it's contribution to the lighting of that fragment and we get a nice result. The only problem is, as the number of lights increases, this naive test becomes much too slow. The solution, pre-process lights into clusters so that we only need to check the lights within a given fragment's cluster. This is the core idea behind the Forward+ approach. -### (TODO: Your README) +# Forward+ +So, what do I mean by clusters. First, consider a small patch of pixels. Just isolating those pixels and the part of the frustrum they span, we get a "mini-frustrum" which see part of our scene. The volume of that mini-frustrum is where we want to check for lights. If a light has a significant effect within that volume, we want to check it with each fragment in that patch. Now, obviously if we were doing a path tracer, then technically all the lights could potentially have an effect since light bounces around all over the scene. But we aren't, so some shortcuts need to be taken since we are prioritizing speed over realism. In this project, our lights were simply point lights. As such, we define a simple radius for each light and that is what determines if a light effects our mini-frustrum. Not we can take this one step furthur. Within this mini-frustrum, there could be a lot of just empty space. This means a basic bounds test would still include a light even though there is no geometry near it to actually do anything. So, we slice that mini-frustrum along the camera Z axis aswell. This gives us small, but precise 3D clusters. If a cluster has geometry in it and a light effects it, then we add that light to that cluster's light list so that our fragments checks it. This takes the number of light check per fragment down significantly and, if implemented correctly, can allow for thousands of lights to be rendered while still keeping good performance. -*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. +Here are some images of what this clustering looks like from the camera: -This assignment has a considerable amount of performance analysis compared -to implementation work. Complete the implementation early to leave time! +![Clusters 64 64 128](IMAGES/CLUSTERS_64_64_128.png) +![Clusters 128 128 128](IMAGES/CLUSTERS_128_128_128.png) +![Clusters 128 128 256](IMAGES/CLUSTERS_128_128_256.png) + +## So what went wrong +If you play around with the live demo, you might notice that Forward+ doesn't really do anything significant. Maybe for a larger number of lights it has slight improvements over naive, but for the most part they are both about the same and both pretty slow. I have an idea as to why, but haven't been able to get the fix working. In my cluster compute shader, I set the light radius that I use in my bounding intersection test to 20 instead of 2. This is intential, because without that, you can't really see anything. However, by setting the radius to be that large, the clusters don't really do anything different from naive because most light intersection most clusters. + +The reason I need to set the radius to 20 has to do with the depth some how. I tried all sorts of ways to linearize the fragment depth to get a nice 0 to 1 depth map, but nothing was working. The best I could do was a hack where you linearize based on a hard coded scene max depth. My compute shader still uses the old depth, which has values only within a small radius around the camera, and thus only light within that region get detected I think. Still not 100% sure, but that is my current guess. + +# Run Time +My Forward+ although it's getting there, is no where near what it needs to be. As such, I don't really have anything to test. The only way for you to be able to see anything is if you turn up the light radius a bunch, but then, as mentioned, this is basically just naive. I could do different cluster sized and configurations, but the results weren't that significant. I couldn't really see any performance differences. Part of this could be my laptop not being that strong of a machine, but overall, knowing the performance results of a buggy implementation I am going to fix didn't seem relevent. ### Credits @@ -30,3 +44,6 @@ to implementation work. Complete the implementation early to leave time! - [dat.GUI](https://github.com/dataarts/dat.gui) - [stats.js](https://github.com/mrdoob/stats.js) - [wgpu-matrix](https://github.com/greggman/wgpu-matrix) +- [Coordinate-Systems](https://learnopengl.com/Getting-started/Coordinate-Systems) +- [depth](https://matthewmacfarquhar.medium.com/webgpu-rendering-part-3-depth-testing-39d4c9ae5bbd) +- CIS 5600 SLIDES diff --git a/src/renderers/forward_plus.ts b/src/renderers/forward_plus.ts index 471796fd..72e0bad8 100644 --- a/src/renderers/forward_plus.ts +++ b/src/renderers/forward_plus.ts @@ -5,16 +5,187 @@ 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 + + // CREATE BIND GROUP "BLUE-PRINT": What is the type of the binding and what can acess that data (doesn't specify the actual data). binding number and visibility is important: MUST MATCH IN SHADER AND BING GROUP INITIALIZATION + this.sceneUniformsBindGroupLayout = renderer.device.createBindGroupLayout({ + label: "scene uniforms bind group layout", + entries: [ + { // CAMERA + binding: 0, + visibility: GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT, + buffer: { type: "uniform"} + }, + { // LIGHTSET + binding: 1, + visibility: GPUShaderStage.FRAGMENT, + buffer: { type: "read-only-storage"} + }, + // { // DEPTH TEXTURE + // binding: 2, + // visibility: GPUShaderStage.FRAGMENT, + // texture: { sampleType: "unfilterable-float" } + // }, + { // CLUSTER PARAMS + binding: 3, + visibility: GPUShaderStage.FRAGMENT, + buffer: { type: "uniform"} + }, + { // CLUSTER LIGHT COUNTS + binding: 4, + visibility: GPUShaderStage.FRAGMENT, + buffer: { type: "storage"} + }, + { // CLUSTER LIGHT INDICES + binding: 5, + visibility: GPUShaderStage.FRAGMENT, + buffer: { type: "storage"} + } + ] + }); + + // ALOCATE MEMORY FOR DEPTH: size of screen + this.depthTexture = renderer.device.createTexture + ({ + size: [renderer.canvas.width, renderer.canvas.height], + format: "depth24plus", + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING + }); + this.depthTextureView = this.depthTexture.createView(); + + // BINDS ACTUAL CPU DATA TO THESE GPU BUFFERS + this.sceneUniformsBindGroup = renderer.device.createBindGroup + ({ + label: "scene uniforms bind group", + layout: this.sceneUniformsBindGroupLayout, + + entries: [ + { // CAMERA + binding: 0, + resource: { buffer: this.camera.uniformsBuffer} + }, + { // LIGHTS + binding: 1, + resource: { buffer: this.lights.lightSetStorageBuffer} + }, + // { // DEPTH + // binding: 2, + // resource: this.linearDepthTextureView + // }, + // // CLUSTER PARAMS + { + binding: 3, + resource: { buffer: this.lights.clusterParamsBuffer} + }, + { // CLUSTER LIGHT COUNT + binding: 4, + resource: { buffer: this.lights.clusterLightCountBuffer} + }, + { // CLUSTER LIGHT IDX + binding: 5, + resource: { buffer: this.lights.clusterLightIdxBuffer} + } + ] + }); + + this.pipeline = renderer.device.createRenderPipeline({ + // BINDS LAYOUTS TO THE PIPELINE? + layout: renderer.device.createPipelineLayout({ + label: "forward+ pipeline layout", + bindGroupLayouts: [ + this.sceneUniformsBindGroupLayout, + renderer.modelBindGroupLayout, + renderer.materialBindGroupLayout + ] + }), + depthStencil: { // + depthWriteEnabled: true, + depthCompare: "less", + format: "depth24plus" + }, + vertex: { // REUSE NAIVE VERTEX SHADER + module: renderer.device.createShaderModule({ + label: "forward+ vert shader", + code: shaders.naiveVertSrc + }), + buffers: [ renderer.vertexBufferLayout ] + }, + fragment: { // USE FORWARD+ FRAG SHADER + module: renderer.device.createShaderModule({ + label: "forward+ 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(); + + // RUN COMPUTE SHADER STUFF + this.lights.populateClusterParamsBuffer(); + this.lights.clearClusterCounts(); + + this.lights.doLightClustering(encoder); + + + // FINAL OUTPUT WHICH WILL USE CLUSTERS FOR LIGHTING (NEEDS UPDATING) + const renderPass = encoder.beginRenderPass({ + label: "naive 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); + 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..98c101f6 100644 --- a/src/renderers/naive.ts +++ b/src/renderers/naive.ts @@ -22,6 +22,11 @@ export class NaiveRenderer extends renderer.Renderer { binding: 1, visibility: GPUShaderStage.FRAGMENT, buffer: { type: "read-only-storage" } + }, + { // CAMERA + binding: 0, + visibility: GPUShaderStage.VERTEX, + buffer: { type: "uniform"} } ] }); @@ -36,6 +41,10 @@ export class NaiveRenderer extends renderer.Renderer { { binding: 1, resource: { buffer: this.lights.lightSetStorageBuffer } + }, + { + binding: 0, + resource: { buffer: this.camera.uniformsBuffer} } ] }); @@ -106,7 +115,8 @@ export class NaiveRenderer extends renderer.Renderer { renderPass.setPipeline(this.pipeline); // TODO-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 => { diff --git a/src/shaders/clustering.cs.wgsl b/src/shaders/clustering.cs.wgsl index 575d6e5a..057f9d06 100644 --- a/src/shaders/clustering.cs.wgsl +++ b/src/shaders/clustering.cs.wgsl @@ -21,3 +21,98 @@ // - Stop adding lights if the maximum number of lights is reached. // - Store the number of lights assigned to this cluster. + +@group(0) @binding(0) var params: ClusterParams; +@group(0) @binding(1) var lightSet: LightSet; +@group(0) @binding(2) var clusterCounts : array; +@group(0) @binding(3) var clusterIndices : array; +@group(0) @binding(4) var uCamera : CameraUniforms; + +fn clusterAABB_view_log(idxX:u32, idxY:u32, d0:f32, d1:f32) -> vec4 { + let tilesX = (params.screenSize.x + params.tileSize.x - 1u) / params.tileSize.x; + let tilesY = (params.screenSize.y + params.tileSize.y - 1u) / params.tileSize.y; + + let xMin = 2.0 * (f32(idxX) / f32(tilesX)) - 1.0; + let xMax = 2.0 * (f32(idxX + 1u) / f32(tilesX)) - 1.0; + let yMin = 2.0 * (f32(idxY) / f32(tilesY)) - 1.0; + let yMax = 2.0 * (f32(idxY + 1u) / f32(tilesY)) - 1.0; + + let tanY = tan(0.5 * params.fovYRadians); + let aspect = f32(params.screenSize.x) / f32(params.screenSize.y); + let Sx = aspect * tanY; + let Sy = tanY; + + let viewX = array( d0*xMin*Sx, d0*xMax*Sx, d1*xMin*Sx, d1*xMax*Sx ); + let viewY = array( d0*yMin*Sy, d0*yMax*Sy, d1*yMin*Sy, d1*yMax*Sy ); + + var xmin = viewX[0]; + var xmax = viewX[0]; + var ymin = viewY[0]; + var ymax = viewY[0]; + + for (var i = 1u; i < 4u; i++) { + xmin = min(xmin, viewX[i]); + xmax = max(xmax, viewX[i]); + ymin = min(ymin, viewY[i]); + ymax = max(ymax, viewY[i]); + } + return vec4(xmin, xmax, ymin, ymax); +} + +fn sphereAABB_intersect(pos:vec3, rad:f32, aabbXY:vec4, zmin:f32, zmax:f32) -> bool { + let qx = clamp(pos.x, aabbXY.x, aabbXY.y); + let qy = clamp(pos.y, aabbXY.z, aabbXY.w); + let qz = clamp(pos.z, zmin, zmax); + let dx = qx - pos.x; + let dy = qy - pos.y; + let dz = qz - pos.z; + return (dx*dx + dy*dy + dz*dz) <= (rad*rad); +} + +@compute +@workgroup_size(4, 4, 4) +fn main(@builtin(global_invocation_id) globalIdx : vec3u) { + let n = params.near; + let f = params.far; + + // cluster dims + let tileX = (params.screenSize.x + params.tileSize.x - 1u) / params.tileSize.x; + let tileY = (params.screenSize.y + params.tileSize.y - 1u) / params.tileSize.y; + + if (globalIdx.x >= tileX || globalIdx.y >= tileY || globalIdx.z >= params.zSlices) { + return; + } + + let idxX = globalIdx.x; + let idxY = globalIdx.y; + let idxZ = globalIdx.z; + + let r = f / n; + let zN = f32(params.zSlices); + let t0 = f32(idxZ) / zN; + let t1 = f32(idxZ + 1u) / zN; + let depthMax = n * pow(r, t0); + let depthMin = n * pow(r, t1); + + let zMin = -depthMin; + let zMax = -depthMax; + + let aabbXY = clusterAABB_view_log(idxX, idxY, depthMax, depthMin); + + let outIdx = (idxZ * tileY + idxY) * tileX + idxX; + + let base = outIdx * params.maxLightsPerCluster; + var count : u32 = 0u; + let lightRad = 20.0; + for (var i = 0u; i < lightSet.numLights; i = i + 1u) { + let P = vec4f(lightSet.lights[i].pos, 1.0); + let Pvs = (uCamera.viewMat * P).xyz; + if (sphereAABB_intersect(Pvs, lightRad, aabbXY, zMin, zMax)) { + if (count < params.maxLightsPerCluster) { + clusterIndices[base + count] = i; + count = count + 1u; + } + } + } + clusterCounts[outIdx] = count; +} \ No newline at end of file diff --git a/src/shaders/common.wgsl b/src/shaders/common.wgsl index 738e9c4e..a672eb72 100644 --- a/src/shaders/common.wgsl +++ b/src/shaders/common.wgsl @@ -12,13 +12,27 @@ struct LightSet { // TODO-2: you may want to create a ClusterSet struct similar to LightSet +struct ClusterParams { + screenSize : vec2u, + tileSize : vec2u, + near : f32, + far : f32, + fovYRadians : f32, + zSlices : u32, + maxLightsPerCluster : u32 +} + struct CameraUniforms { // TODO-1.3: add an entry for the view proj mat (of type mat4x4f) + viewProjMat: mat4x4f, + viewMat: mat4x4f, + nearFar: vec2f } // 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); + // HARD CODED 2 + return clamp(1.f - pow(distance / 2.0, 4.f), 0.f, 1.f) / (distance * distance); } fn calculateLightContrib(light: Light, posWorld: vec3f, nor: vec3f) -> vec3f { diff --git a/src/shaders/forward_plus.fs.wgsl b/src/shaders/forward_plus.fs.wgsl index 0500e3df..5c066821 100644 --- a/src/shaders/forward_plus.fs.wgsl +++ b/src/shaders/forward_plus.fs.wgsl @@ -14,3 +14,104 @@ // 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(0) @binding(0) var uCamera : CameraUniforms; +@group(0) @binding(1) var lightSet: LightSet; +@group(0) @binding(3) var params: ClusterParams; +@group(0) @binding(4) var clusterCounts: array; +@group(0) @binding(5) var clusterIdx: array; + +@group(2) @binding(0) var diffuseTex: texture_2d; +@group(2) @binding(1) var diffuseTexSampler: sampler; + +struct FragmentInput +{ + @builtin(position) fragCoord: vec4f, + @location(0) pos: vec3f, + @location(1) nor: vec3f, + @location(2) uv: vec2f, + @location(3) nearPlane: f32, + @location(4) farPlane: f32 +} + +fn hash3(u: u32) -> vec3f { + let x = f32(((u * 1664525u) ^ 1013904223u) & 1023u) / 1023.0; + let y = f32(((u * 22695477u) ^ 1u) & 1023u) / 1023.0; + let z = f32(((u * 1103515245u) ^ 12345u) & 1023u) / 1023.0; + return vec3f(x, y, z); +} + +fn tilesXY() -> vec2 { + let tx = (params.screenSize.x + params.tileSize.x - 1u) / params.tileSize.x; + let ty = (params.screenSize.y + params.tileSize.y - 1u) / params.tileSize.y; + return vec2(tx, ty); +} + +fn clusterId(ix:u32, iy:u32, iz:u32) -> u32 { + let t = tilesXY(); + return (iz * t.y + iy) * t.x + ix; +} + +fn sliceZ_fromDepth_log(d: f32, n: f32, f: f32, zSlices: u32) -> u32 { + let t = log(clamp(d, n, f) / n) / log(f / n); + return u32(floor(clamp(t, 0.0, 0.999999) * f32(zSlices))); +} + +@fragment +fn main(in: FragmentInput) -> @location(0) vec4f +{ + let n = in.nearPlane; + let f = in.farPlane; + + // RENDER CLUSTERS WITH HASH COLORS + // let tileX = u32(in.fragCoord.x) / params.tileSize.x; + // let tileY = u32(in.fragCoord.y) / params.tileSize.y; + + //let dwow = in.fragCoord.z; + // let z_ndc = dwow * 2.0 - 1.0; + // let linearDepth = (2.0 * n * f) / (f + n - z_ndc * (f - n)); + // let linear = clamp(linearDepth / 100.0, 0.0, 1.0); + + // let z = clamp(linear, n, f); + // let dist = log(z / n) / log(f / n); + // let slice = u32(floor(dist * f32(params.zSlices))); + // let magic = clamp(slice, 0, params.zSlices - 1u); + + // return vec4(hash3(tileX + tileY + magic), 1.0); + + // Alpha test + let albedo = textureSample(diffuseTex, diffuseTexSampler, in.uv); + if (albedo.a < 0.5) { + discard; + } + + let tilesX = (params.screenSize.x + params.tileSize.x - 1u) / params.tileSize.x; + let tilesY = (params.screenSize.y + params.tileSize.y - 1u) / params.tileSize.y; + let indexX = min(u32(in.fragCoord.x) / params.tileSize.x, tilesX - 1u); + let indexY = min(u32(in.fragCoord.y) / params.tileSize.y, tilesY - 1u); + + // THIS IS KIND OF JANK, BUT IT"S THE ONLY THING THAT SEEMED TO GIVE ANY SEMI-ACCURATE [0,1] DEPTH RANGE VALUES. + let dwow = in.fragCoord.z; + let z_ndc = dwow * 2.0 - 1.0; + let linearDepth = (2.0 * n * f) / (f + n - z_ndc * (f - n)); + let depth = clamp(linearDepth / 100.0, 0.0, 1.0); // CHANGE 100.0 based on scene depth + + let logZIndex = sliceZ_fromDepth_log(depth, n, f, params.zSlices); + + let clusterIndex = clusterId(indexX, indexY, logZIndex); + + let base = clusterIndex * params.maxLightsPerCluster; + let count = clusterCounts[clusterIndex]; + + let N = normalize(in.nor); + var total = vec3f(0.0, 0.0, 0.0); + + for (var k:u32 = 0u; k < count; k = k + 1u) { + let li = clusterIdx[base + k]; + if (li < lightSet.numLights) { + let L = lightSet.lights[li]; + total += calculateLightContrib(L, in.pos, normalize(in.nor)); + } + } + + return vec4f(albedo.rgb * total, 1.0); +} diff --git a/src/shaders/move_lights.cs.wgsl b/src/shaders/move_lights.cs.wgsl index 89c05f5e..5b0bb5bb 100644 --- a/src/shaders/move_lights.cs.wgsl +++ b/src/shaders/move_lights.cs.wgsl @@ -1,5 +1,5 @@ -@group(${bindGroup_scene}) @binding(0) var lightSet: LightSet; -@group(${bindGroup_scene}) @binding(1) var time: f32; +@group(0) @binding(0) var lightSet: LightSet; +@group(0) @binding(1) var time: f32; // https://gist.github.com/munrocket/236ed5ba7e409b8bdf1ff6eca5dcdc39 // MIT License. © Stefan Gustavson, Munrocket diff --git a/src/shaders/naive.fs.wgsl b/src/shaders/naive.fs.wgsl index 0afeaac7..09758024 100644 --- a/src/shaders/naive.fs.wgsl +++ b/src/shaders/naive.fs.wgsl @@ -1,13 +1,13 @@ -@group(${bindGroup_scene}) @binding(1) var lightSet: LightSet; +@group(0) @binding(1) var lightSet: LightSet; -@group(${bindGroup_material}) @binding(0) var diffuseTex: texture_2d; -@group(${bindGroup_material}) @binding(1) var diffuseTexSampler: sampler; +@group(2) @binding(0) var diffuseTex: texture_2d; +@group(2) @binding(1) var diffuseTexSampler: sampler; struct FragmentInput { @location(0) pos: vec3f, @location(1) nor: vec3f, - @location(2) uv: vec2f + @location(2) uv: vec2f, } @fragment diff --git a/src/shaders/naive.vs.wgsl b/src/shaders/naive.vs.wgsl index 5a7ddd4b..5c841a36 100644 --- a/src/shaders/naive.vs.wgsl +++ b/src/shaders/naive.vs.wgsl @@ -3,7 +3,9 @@ // TODO-1.3: add a uniform variable here for camera uniforms (of type CameraUniforms) // make sure to use ${bindGroup_scene} for the group -@group(${bindGroup_model}) @binding(0) var modelMat: mat4x4f; + +@group(0) @binding(0) var uCamera : CameraUniforms; +@group(1) @binding(0) var modelMat: mat4x4f; struct VertexInput { @@ -14,10 +16,12 @@ struct VertexInput struct VertexOutput { - @builtin(position) fragPos: vec4f, + @builtin(position) fragCoord: vec4f, @location(0) pos: vec3f, @location(1) nor: vec3f, - @location(2) uv: vec2f + @location(2) uv: vec2f, + @location(3) nearPlane: f32, + @location(4) farPlane: f32 } @vertex @@ -26,9 +30,15 @@ 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.fragCoord = uCamera.viewProjMat * modelPos; // TODO-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; + + let test = uCamera.viewMat * modelPos; + + out.nearPlane = uCamera.nearFar.x; + out.farPlane = uCamera.nearFar.y; + return out; } diff --git a/src/shaders/shaders.ts b/src/shaders/shaders.ts index 584c008f..7cfadb7c 100644 --- a/src/shaders/shaders.ts +++ b/src/shaders/shaders.ts @@ -29,20 +29,37 @@ export const constants = { bindGroup_material: 2, moveLightsWorkgroupSize: 128, + clusterWorkgroupSize: 1, lightRadius: 2 }; // ================================= -function evalShaderRaw(raw: string) { - return eval('`' + raw.replaceAll('${', '${constants.') + '`'); +// function evalShaderRaw(raw: string) { +// return eval('`' + raw.replaceAll('${', '${constants.') + '`'); +// } + +// const commonSrc: string = evalShaderRaw(commonRaw); + +// function processShaderRaw(raw: string) { +// return commonSrc + evalShaderRaw(raw); +// } + +const RE_CONST = /\$\{\s*(?:constants\.)?([A-Za-z_]\w*)\s*\}/g; +function substituteConstants(raw: string): string { + return raw.replace(RE_CONST, (_m, key) => { + const val = (constants as Record)[key]; + if (val === undefined) { + throw new Error(`Shader constant "${key}" not found in constants`); + } + return String(val); + }); } -const commonSrc: string = evalShaderRaw(commonRaw); - +const commonSrc: string = substituteConstants(commonRaw); function processShaderRaw(raw: string) { - return commonSrc + evalShaderRaw(raw); + return commonSrc + substituteConstants(raw); } export const naiveVertSrc: string = processShaderRaw(naiveVertRaw); diff --git a/src/stage/camera.ts b/src/stage/camera.ts index 7d2a4a1e..5fb4fe2c 100644 --- a/src/stage/camera.ts +++ b/src/stage/camera.ts @@ -1,16 +1,33 @@ -import { Mat4, mat4, Vec3, vec3 } from "wgpu-matrix"; +import { Mat4, mat4, vec2, 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); + readonly buffer = new ArrayBuffer(36 * 4); private readonly floatView = new Float32Array(this.buffer); set viewProjMat(mat: Float32Array) { // TODO-1.1: set the first 16 elements of `this.floatView` to the input `mat` + for (let index = 0; index < this.floatView.length; index++) { + this.floatView[index] = mat[index]; + } } + private readonly view = new Float32Array(this.buffer); + set viewMat(mat: Float32Array) { + for (let index = 0; index < this.view.length; index++) { + this.view[index+16] = mat[index]; + } + } + + private readonly f32 = new Float32Array(this.buffer); + set nearFar(nf: Float32Array) { + this.f32[32] = nf[0]; // near + this.f32[33] = nf[1]; // far + } + // TODO-2: add extra functions to set values needed for light clustering here + } export class Camera { @@ -39,6 +56,12 @@ export class Camera { // // 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: this.uniforms.buffer.byteLength, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST + }); + this.projMat = mat4.perspective(toRadians(fovYDegrees), aspectRatio, Camera.nearPlane, Camera.farPlane); this.rotateCamera(0, 0); // set initial camera vectors @@ -130,9 +153,17 @@ export class Camera { const viewProjMat = mat4.mul(this.projMat, viewMat); // TODO-1.1: set `this.uniforms.viewProjMat` to the newly calculated view proj mat + this.uniforms.viewProjMat = viewProjMat; + this.uniforms.viewMat = viewMat; + + const nfPlanes = vec2.create(Camera.nearPlane, Camera.farPlane); + this.uniforms.nearFar = nfPlanes; + // TODO-2: write to extra buffers needed for light clustering here // TODO-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, this.uniforms.buffer); } } diff --git a/src/stage/lights.ts b/src/stage/lights.ts index a6eed919..db01ebc1 100644 --- a/src/stage/lights.ts +++ b/src/stage/lights.ts @@ -3,6 +3,7 @@ import { device } from "../renderer"; import * as shaders from '../shaders/shaders'; import { Camera } from "./camera"; +import { canvas, fovYDegrees } from "../renderer" // GET SCREEN ZISE // h in [0, 1] function hueToRgb(h: number) { @@ -13,7 +14,7 @@ function hueToRgb(h: number) { export class Lights { private camera: Camera; - numLights = 500; + numLights = 100; // STARTING WITH 100 static readonly maxNumLights = 5000; static readonly numFloatsPerLight = 8; // vec3f is aligned at 16 byte boundaries @@ -30,6 +31,24 @@ export class Lights { // TODO-2: add layouts, pipelines, textures, etc. needed for light clustering here + // DEFAULT CLUSTER PARAMS: + clusterWidth = 128; + clusterHeight = 128; + zSlice = 256; + maxLightPerCluster = 500; + + tileX = 0; + tileY = 0; + numClusters = 0; + + clusterParamsBuffer: GPUBuffer; + clusterLightCountBuffer: GPUBuffer; + clusterLightIdxBuffer: GPUBuffer; + + clusterBindGroupLayout: GPUBindGroupLayout; + clusterBindGroup: GPUBindGroup; + clusterComputePipeline: GPUComputePipeline; + constructor(camera: Camera) { this.camera = camera; @@ -94,6 +113,129 @@ export class Lights { }); // TODO-2: initialize layouts, pipelines, textures, etc. needed for light clustering here + + this.tileX = Math.ceil(canvas.width / this.clusterWidth); + this.tileY = Math.ceil(canvas.height / this.clusterHeight); + this.numClusters = this.tileX * this.tileY * this.zSlice; + + this.clusterParamsBuffer = device.createBuffer({ + label: "cluster params", + size: 64, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST + }); + + this.clusterLightCountBuffer = device.createBuffer({ + label: "cluster light count", + size: this.numClusters * 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + }); + + this.clusterLightIdxBuffer = device.createBuffer({ + label: "cluster light idxs", + size: this.numClusters * this.maxLightPerCluster * 4, + usage: GPUBufferUsage.STORAGE + }); + + this.populateClusterParamsBuffer(); + + this.clusterBindGroupLayout = device.createBindGroupLayout({ + label: "bind cluster group layout", + entries: [ + { // CLUSTER PARAMS + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: "uniform"} + }, + { // LIGHTS + binding: 1, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: "read-only-storage"} + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: "storage"} + }, + { + binding: 3, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: "storage"} + }, + { + binding: 4, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: "uniform"} + } + ] + }); + + this.clusterBindGroup = device.createBindGroup({ + label: "bind cluster bind group", + layout: this.clusterBindGroupLayout, + entries: [ + { + binding: 0, + resource: { buffer: this.clusterParamsBuffer} + }, + { + binding: 1, + resource: { buffer: this.lightSetStorageBuffer} + }, + { + binding: 2, + resource: { buffer: this.clusterLightCountBuffer} + }, + { + binding: 3, + resource: { buffer: this.clusterLightIdxBuffer} + }, + { + binding: 4, + resource: { buffer: this.camera.uniformsBuffer} + } + ] + }); + + this.clusterComputePipeline = device.createComputePipeline({ + label: "cluster compute pipeline", + layout: device.createPipelineLayout({ + label: "cluster compute pipeline layout", + bindGroupLayouts: [ this.clusterBindGroupLayout] + }), + compute: { + module: device.createShaderModule({ + label: "cluster compute shader", + code: shaders.clusteringComputeSrc + }), + entryPoint: "main" + } + }); + } + + populateClusterParamsBuffer() { + const u = new Uint32Array(12); // 12 * 4 = 48 bytes + const f = new Float32Array(u.buffer); + const fovYRad = fovYDegrees * (Math.PI / 180); + + u[0] = canvas.width; + u[1] = canvas.height; + u[2] = this.clusterWidth; + u[3] = this.clusterHeight; + + f[4] = Camera.nearPlane; // byte 16 + f[5] = Camera.farPlane; // byte 20 + f[6] = fovYRad; // byte 24 + u[7] = this.zSlice; // byte 28 + u[8] = this.maxLightPerCluster; // byte 32 + // bytes 36..47 padding (u[9..11] unused) + + device.queue.writeBuffer(this.clusterParamsBuffer, 0, u); + } + + clearClusterCounts() { + device.queue.writeBuffer( + this.clusterLightCountBuffer, 0, new Uint8Array(this.numClusters * 4) + ); } private populateLightsBuffer() { @@ -113,6 +255,20 @@ export class Lights { doLightClustering(encoder: GPUCommandEncoder) { // TODO-2: run the light clustering compute pass(es) here // implementing clustering here allows for reusing the code in both Forward+ and Clustered Deferred + const tileX = Math.ceil(canvas.width / this.clusterWidth); + const tileY = Math.ceil(canvas.height / this.clusterHeight); + + const pass = encoder.beginComputePass({ label: "cluster stub" }); + pass.setPipeline(this.clusterComputePipeline); + pass.setBindGroup(0, this.clusterBindGroup); + + const wgx = 4, wgy = 4, wgz = 4; + const gx = Math.ceil(tileX / wgx); + const gy = Math.ceil(tileY / wgy); + const gz = Math.ceil(this.zSlice / wgz); + + pass.dispatchWorkgroups(gx, gy, gz); + pass.end(); } // CHECKITOUT: this is where the light movement compute shader is dispatched from the host