diff --git a/README.md b/README.md index f99cdff..0b425d3 100644 --- a/README.md +++ b/README.md @@ -1,31 +1,55 @@ -# Project5-WebGPU-Gaussian-Splat-Viewer +# WebGPU Gaussian Splat Viewer -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 4** +**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 5** -* (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) +* Mufeng Xu +* Tested on: **Google Chrome 129.0** on + Windows 11, i9-13900H @ 2.6GHz 32GB, RTX 4080 Laptop 12282MB (Personal Computer) -### Live Demo +## Live Demo -[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +[Live Demo](https://solemnwind.github.io/Project5-WebGPU-Gaussian-Splat-Viewer/) (Open with Chrome) -### Demo Video/GIF +## Demo Video/GIF -[![](img/video.mp4)](TODO) +![](images/bonsai.gif) -### (TODO: Your README) +## Project Overview -*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. +Gaussian Splatting is a volume rendering technique that deals with the direct rendering of volume data without converting the data into surface or line primitives. This method uses 3D gaussian distributions to represent objects, with spherical harmonics, the gaussian splats can represent sophisticated colors and details. This project specifically implements a viewer that renders a gaussian splat scene. -This assignment has a considerable amount of performance analysis compared -to implementation work. Complete the implementation early to leave time! +**Implemented Functions** -### Credits +* Add MVP calculation to Point Cloud rendering +* View frustum culling to remove non-visible splats +* Use spherical harmonics to evaluate colors of splats +* Gaussian Splatter renderer -- [Vite](https://vitejs.dev/) -- [tweakpane](https://tweakpane.github.io/docs//v3/monitor-bindings/) -- [stats.js](https://github.com/mrdoob/stats.js) -- [wgpu-matrix](https://github.com/greggman/wgpu-matrix) -- Special Thanks to: Shrek Shao (Google WebGPU team) & [Differential Guassian Renderer](https://github.com/graphdeco-inria/diff-gaussian-rasterization) +## Performance Analysis + +### Compare your results from point-cloud and gaussian renderer, what are the differences? + +The point-cloud renderer shows sparse points, these points typically don't have volumes and can't block each other, +and they don't have transparency and color (but color is possible, though not detailed). +While the gaussian renderer renders globs at the positions of the points, the scene is smoother and much more realistic. +Because the use of spherical harmonics for representing the colors, gaussian splatting renders detailed colors with a rather low cost. + +### For gaussian renderer, how does changing the workgroup-size affect performance? Why do you think this is? + +Changing the workgroup size larger, generally makes the GPU more utilized, therefore improve the performance. However, if the workgroup is too large, first the increased data contention will slow down the performance, second, for some scenes, an increased workgroup causes insufficient buffer space, and the renderer does not work in that case. While decreasing the workgroup size will lead to less parallel computing and lower the FPS. + +### Does view-frustum culling give performance improvement? Why do you think this is? + +In the cases where the camera is *inside* the scene, such as the room, view frustum culling could lower the GPU usage, and potentially improve the FPS if the scene is large. And this is also what I observed: even for a simple scene, why zoomed out and include the whole scene in the screen, the GPU usage rises immediately, and for a larger scene, it causes lower FPS. However, view-frustum culling increases overhead, potentially worsen the improvement in simple scenes or the case that most of the gaussians are in the view-frustum. + +### Does number of gaussians affect performance? Why do you think this is? + +When the number is not too large, because of the parallelism, it doesn't impact the performance. In my case, before the buffer space ran out, the FPS was always steady at ~120. But the increased number of gaussians makes culling, sorting etc. more expensive, and therefore impact the performance. + +## Credits + +* [Vite](https://vitejs.dev/) +* [tweakpane](https://tweakpane.github.io/docs//v3/monitor-bindings/) +* [stats.js](https://github.com/mrdoob/stats.js) +* [wgpu-matrix](https://github.com/greggman/wgpu-matrix) +* Special Thanks to: Shrek Shao (Google WebGPU team) & [Differential Guassian Renderer](https://github.com/graphdeco-inria/diff-gaussian-rasterization) diff --git a/images/bonsai.gif b/images/bonsai.gif new file mode 100644 index 0000000..b8dbc1d Binary files /dev/null and b/images/bonsai.gif differ diff --git a/src/renderers/gaussian-renderer.ts b/src/renderers/gaussian-renderer.ts index 1684523..b34e403 100644 --- a/src/renderers/gaussian-renderer.ts +++ b/src/renderers/gaussian-renderer.ts @@ -5,7 +5,7 @@ import { get_sorter,c_histogram_block_rows,C } from '../sort/sort'; import { Renderer } from './renderer'; export interface GaussianRenderer extends Renderer { - + render_setting_buffer: GPUBuffer } // Utility to create GPU buffers @@ -35,6 +35,31 @@ export default function get_renderer( // =============================================== const nulling_data = new Uint32Array([0]); + const nulling_buffer = createBuffer( + device, + 'nulling buffer', + 4, + GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, + nulling_data + ); + + const render_setting_buffer = createBuffer( + device, + 'render setting buffer', + 8, + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + new Float32Array([ + 1.0, + pc.sh_deg + ]) + ); + + const splat_buffer = createBuffer( + device, + 'splat buffer', + pc.num_points * 24 * 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); // =============================================== // Create Compute Pipeline and Bind Groups @@ -52,6 +77,25 @@ export default function get_renderer( }, }); + const uniform_bind_group = device.createBindGroup({ + label: 'uniform_bind_group', + layout: preprocess_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer }}, + { binding: 1, resource: { buffer: render_setting_buffer }}, + ], + }); + + const gaussian_bind_group = device.createBindGroup({ + label: 'gaussian_bind_group', + layout: preprocess_pipeline.getBindGroupLayout(1), + entries: [ + { binding: 0, resource: { buffer: pc.gaussian_3d_buffer }}, + { binding: 1, resource: { buffer: splat_buffer}}, + { binding: 2, resource: { buffer: pc.sh_buffer}}, + ], + }); + const sort_bind_group = device.createBindGroup({ label: 'sort', layout: preprocess_pipeline.getBindGroupLayout(2), @@ -68,19 +112,111 @@ export default function get_renderer( // Create Render Pipeline and Bind Groups // =============================================== + // create the indirect buffer + const indirect_buffer = createBuffer( + device, + 'indirect buffer', + 16, + GPUBufferUsage.COPY_DST | GPUBufferUsage.INDIRECT, + new Uint32Array([ 6, 0, 0, 0 ]) + ); + + // Render pipeline + const render_pipeline = device.createRenderPipeline({ + label: 'render pipeline', + layout: 'auto', + vertex: { + module: device.createShaderModule({ + code: renderWGSL + }), + entryPoint: 'vs_main' + }, + fragment: { + module: device.createShaderModule({ + code: renderWGSL + }), + targets: [{ + format: presentation_format, + blend: { + color: { + srcFactor: 'one', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + alpha: { + srcFactor: 'one', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + } + }], + entryPoint: 'fs_main' + } + }); + + const render_bind_group = device.createBindGroup({ + label: 'render bind group', + layout: render_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer }}, + { binding: 1, resource: { buffer: splat_buffer }}, + { binding: 2, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer }}, + ] + }); + // =============================================== // Command Encoder Functions // =============================================== - + const preprocess = (encoder: GPUCommandEncoder) => { + const preprocess_pass = encoder.beginComputePass(); + preprocess_pass.setPipeline(preprocess_pipeline); + preprocess_pass.setBindGroup(0, uniform_bind_group); + preprocess_pass.setBindGroup(1, gaussian_bind_group); + preprocess_pass.setBindGroup(2, sort_bind_group); + preprocess_pass.dispatchWorkgroups(Math.ceil(pc.num_points / C.histogram_wg_size)); + preprocess_pass.end(); + }; // =============================================== // Return Render Object // =============================================== return { frame: (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + encoder.copyBufferToBuffer( + nulling_buffer, 0, + sorter.sort_info_buffer, 0, + 4); + encoder.copyBufferToBuffer( + nulling_buffer, 0, + sorter.sort_dispatch_indirect_buffer, 0, + 4); + preprocess(encoder); sorter.sort(encoder); + encoder.copyBufferToBuffer( + sorter.sort_info_buffer, 0, + indirect_buffer, 4, + 4 + ); + + + // Render pass + const render_pass = encoder.beginRenderPass({ + label: 'gaussian render pass', + colorAttachments: [{ + view: texture_view, + loadOp: 'clear', + storeOp: 'store' + }] + }); + + render_pass.setPipeline(render_pipeline); + console.log("setBindGroup"); + render_pass.setBindGroup(0, render_bind_group); + render_pass.drawIndirect(indirect_buffer, 0); + render_pass.end(); }, camera_buffer, + render_setting_buffer }; } diff --git a/src/renderers/renderer.ts b/src/renderers/renderer.ts index ffdf9ba..d5c78bb 100644 --- a/src/renderers/renderer.ts +++ b/src/renderers/renderer.ts @@ -122,6 +122,13 @@ export default async function init( {min: 0, max: 1.5} ).on('change', (e) => { //TODO: Bind constants to the gaussian renderer. + if (gaussian_renderer) { + device.queue.writeBuffer( + gaussian_renderer.render_setting_buffer, + 0, + new Float32Array([e.value]) + ); + } }); } diff --git a/src/shaders/gaussian.wgsl b/src/shaders/gaussian.wgsl index 759226d..fbdd2f9 100644 --- a/src/shaders/gaussian.wgsl +++ b/src/shaders/gaussian.wgsl @@ -1,22 +1,80 @@ struct VertexOutput { @builtin(position) position: vec4, //TODO: information passed from vertex shader to fragment shader + @location(0) center: vec2f, + @location(1) color: vec4f, + @location(2) conic: vec3f, + @location(3) opacity: f32 +}; + +struct CameraUniforms { + view: mat4x4, + view_inv: mat4x4, + proj: mat4x4, + proj_inv: mat4x4, + viewport: vec2, + focal: vec2 }; struct Splat { //TODO: information defined in preprocess compute shader + packed_pos: u32, + packed_size: u32, + packed_color: array, + packed_conic_opacity: array }; +@group(0) @binding(0) var camera: CameraUniforms; +@group(0) @binding(1) var splat_buffer: array; +@group(0) @binding(2) var sorting_info: array; + @vertex fn vs_main( + @builtin(instance_index) instance_idx: u32, + @builtin(vertex_index) vertex_idx: u32 ) -> VertexOutput { //TODO: reconstruct 2D quad based on information from splat, pass var out: VertexOutput; - out.position = vec4(1. ,1. , 0., 1.); + + let splat_idx = sorting_info[instance_idx]; + let splat = splat_buffer[splat_idx]; + + let pos = unpack2x16float(splat.packed_pos); + let size = unpack2x16float(splat.packed_size); + let x = pos.x; + let y = pos.y; + let w = size.x; + let h = size.y; + + let vertices = array( + vec2f(x - w, y + h), + vec2f(x - w, y - h), + vec2f(x + w, y - h), + vec2f(x + w, y - h), + vec2f(x + w, y + h), + vec2f(x - w, y + h) + ); + out.position = vec4f(vertices[vertex_idx], 0.0, 1.0); + + out.color = vec4f(unpack2x16float(splat.packed_color[0]), + unpack2x16float(splat.packed_color[1])); + out.conic = vec3f(unpack2x16float(splat.packed_conic_opacity[0]), + unpack2x16float(splat.packed_conic_opacity[1]).x); + out.opacity = unpack2x16float(splat.packed_conic_opacity[1]).y; + out.center = (0.5 + pos * vec2f(0.5, -0.5)) * camera.viewport; + return out; } @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4 { - return vec4(1.); + // https://github.com/graphdeco-inria/diff-gaussian-rasterization/blob/main/cuda_rasterizer/forward.cu#L263 + let d = in.center - in.position.xy; + let p = -0.5 * (in.conic.x * d.x * d.x + in.conic.z * d.y * d.y) + - in.conic.y * d.x * d.y; + + if (p > 0.0) { return vec4f(0.0, 0.0, 0.0, 0.0); } + + let alpha = min(0.99f, in.opacity * exp(p)); + return in.color * alpha; } \ No newline at end of file diff --git a/src/shaders/point_cloud.wgsl b/src/shaders/point_cloud.wgsl index 01dded1..617171e 100644 --- a/src/shaders/point_cloud.wgsl +++ b/src/shaders/point_cloud.wgsl @@ -35,7 +35,7 @@ fn vs_main( let pos = vec4(a.x, a.y, b.x, 1.); // TODO: MVP calculations - out.position = pos; + out.position = camera.proj * camera.view * pos; return out; } diff --git a/src/shaders/preprocess.wgsl b/src/shaders/preprocess.wgsl index bbc63f5..41a11e1 100644 --- a/src/shaders/preprocess.wgsl +++ b/src/shaders/preprocess.wgsl @@ -24,7 +24,7 @@ struct DispatchIndirect { dispatch_x: atomic, dispatch_y: u32, dispatch_z: u32, -} +}; struct SortInfos { keys_size: atomic, // instance_count in DrawIndirect @@ -33,7 +33,7 @@ struct SortInfos { passes: u32, even_pass: u32, odd_pass: u32, -} +}; struct CameraUniforms { view: mat4x4, @@ -47,7 +47,7 @@ struct CameraUniforms { struct RenderSettings { gaussian_scaling: f32, sh_deg: f32, -} +}; struct Gaussian { pos_opacity: array, @@ -56,10 +56,25 @@ struct Gaussian { }; struct Splat { - //TODO: store information for 2D splat rendering + packed_pos: u32, + packed_size: u32, + packed_color: array, + packed_conic_opacity: array, }; //TODO: bind your data here +@group(0) @binding(0) +var camera: CameraUniforms; +@group(0) @binding(1) +var render_settings: RenderSettings; + +@group(1) @binding(0) +var gaussians: array; +@group(1) @binding(1) +var splats: array; +@group(1) @binding(2) +var colors: array; + @group(2) @binding(0) var sort_infos: SortInfos; @group(2) @binding(1) @@ -72,7 +87,15 @@ var sort_dispatch: DispatchIndirect; /// reads the ith sh coef from the storage buffer fn sh_coef(splat_idx: u32, c_idx: u32) -> vec3 { //TODO: access your binded sh_coeff, see load.ts for how it is stored - return vec3(0.0); + let index = splat_idx * 24 + c_idx % 2 + c_idx / 2 * 3; + let color_a_b = unpack2x16float(colors[index]); + let color_c_d = unpack2x16float(colors[index + 1]); + + if (c_idx % 2 == 0) { + return vec3f(color_a_b.x, color_a_b.y, color_c_d.x); + } else { + return vec3f(color_a_b.y, color_c_d.x, color_c_d.y); + } } // spherical harmonics evaluation with Condon–Shortley phase @@ -111,8 +134,111 @@ fn computeColorFromSH(dir: vec3, v_idx: u32, sh_deg: u32) -> vec3 { @compute @workgroup_size(workgroupSize,1,1) fn preprocess(@builtin(global_invocation_id) gid: vec3, @builtin(num_workgroups) wgs: vec3) { let idx = gid.x; - //TODO: set up pipeline as described in instruction + if (idx >= arrayLength(&gaussians)) { + return; + } + + // extract gaussian information + let gaussian = gaussians[idx]; + let pos_xy = unpack2x16float(gaussian.pos_opacity[0]); + let pos_za = unpack2x16float(gaussian.pos_opacity[1]); + let pos = vec4f(pos_xy, pos_za.x, 1.0); + let opacity = 1.0f / (1.0f + exp(-pos_za.y)); + + // get ndc + let view_pos = camera.view * pos; + var ndc = camera.proj * view_pos; + ndc /= ndc.w; + + // view-frustum culling + if (ndc.x < -1.2 || ndc.x > 1.2 || + ndc.y < -1.2 || ndc.y > 1.2 || + view_pos.z <= 0.0) { + return; + } + + // get rotation + let rot_wx = unpack2x16float(gaussian.rot[0]); + let rot_yz = unpack2x16float(gaussian.rot[1]); + let w = rot_wx.x; + let x = rot_wx.y; + let y = rot_yz.x; + let z = rot_yz.y; + + let R = mat3x3f( + 1.0 - 2.0 * (y * y + z * z), 2.0 * (x * y - w * z) , 2.0 * (x * z + w * y), + 2.0 * (x * y + w * z) , 1.0 - 2.0 * (x * x + z * z), 2.0 * (y * z - w * x), + 2.0 * (x * z - w * y) , 2.0 * (y * z + w * x) , 1.0 - 2.0 * (x * x + y * y) + ); + + // get scale + let scale_xy = exp(unpack2x16float(gaussian.scale[0])); + let scale_zw = exp(unpack2x16float(gaussian.scale[1])); + let sx = scale_xy.x * render_settings.gaussian_scaling; + let sy = scale_xy.y * render_settings.gaussian_scaling; + let sz = scale_zw.x * render_settings.gaussian_scaling; + let S = mat3x3f( + sx, 0.0, 0.0, + 0.0, sy, 0.0, + 0.0, 0.0, sz + ); + + let Cov3D = transpose(R) * transpose(S) * S * R; + + // Jacobian + let J = mat3x3f( + camera.focal.x / view_pos.z, 0.0, -camera.focal.x * view_pos.x / (view_pos.z * view_pos.z), + 0.0, camera.focal.y / view_pos.z, -camera.focal.y * view_pos.y / (view_pos.z * view_pos.z), + 0.0, 0.0, 0.0 + ); + + let W = transpose(mat3x3f( + camera.view[0].xyz, camera.view[1].xyz, camera.view[2].xyz + )); + + let WJ = W * J; + + let Sigma = mat3x3f( + Cov3D[0][0], Cov3D[0][1], Cov3D[0][2], + Cov3D[0][1], Cov3D[1][1], Cov3D[1][2], + Cov3D[0][2], Cov3D[1][2], Cov3D[2][2] + ); + + var Cov2D = transpose(WJ) * Sigma * WJ; + Cov2D[0][0] += 0.3; + Cov2D[1][1] += 0.3; + let cxx = Cov2D[0][0]; + let cyy = Cov2D[1][1]; + let cxy = Cov2D[0][1]; + + let det = cxx * cyy - cxy * cxy; + if (det == 0.0) { return; } + + let mid = (cxx + cyy) * 0.5; + let lambda1 = mid + sqrt(max(0.1, mid * mid - det)); + let lambda2 = mid - sqrt(max(0.1, mid * mid - det)); + let radius = ceil(3.0f * sqrt(max(lambda1, lambda2))); + + let cam_pos = -camera.view[3].xyz; + let direction = normalize(pos.xyz- cam_pos); + let color = computeColorFromSH(direction, idx, u32(render_settings.sh_deg)); + + let conic = vec3f(cyy / det, -cxy / det, cxx / det); + + let sorted_idx = atomicAdd(&sort_infos.keys_size, 1); + splats[sorted_idx].packed_pos = pack2x16float(ndc.xy); + splats[sorted_idx].packed_size = pack2x16float(vec2f(radius, radius) / camera.viewport); + splats[sorted_idx].packed_color[0] = pack2x16float(color.rg); + splats[sorted_idx].packed_color[1] = pack2x16float(vec2f(color.b, 1.0f)); + splats[sorted_idx].packed_conic_opacity[0] = pack2x16float(conic.xy); + splats[sorted_idx].packed_conic_opacity[1] = pack2x16float(vec2f(conic.z, opacity)); + + sort_depths[sorted_idx] = bitcast(100.0 - view_pos.z); + sort_indices[sorted_idx] = sorted_idx; let keys_per_dispatch = workgroupSize * sortKeyPerThread; // increment DispatchIndirect.dispatchx each time you reach limit for one dispatch of keys + if (sorted_idx % keys_per_dispatch == 0) { + atomicAdd(&sort_dispatch.dispatch_x, 1); + } } \ No newline at end of file