diff --git a/README.md b/README.md index edffdaf..c8fbb0e 100644 --- a/README.md +++ b/README.md @@ -2,25 +2,41 @@ **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) +* Christina Qiu + * [LinkedIn](https://www.linkedin.com/in/christina-qiu-6094301b6/), [personal website](https://christinaqiu3.github.io/), [twitter](), etc. +* Tested on: Windows 11, Intel Core i7-13700H @ 2.40GHz, 16GB RAM, NVIDIA GeForce RTX 4060 Laptop GPU (Personal laptop) + +## Overview + +This project implements a GPU-driven 3D Gaussian splatting renderer using WebGPU. The goal is to take a set of 3D Gaussians and render high-quality, order-independent, depth-sorted 2D splats (quads) on the screen. The implementation follows the assignment stages: loading scene and camera data, preprocessing Gaussians on the GPU (culling, covariance → 2D conic projection, color evaluation), sorting by depth, and rendering the resulting splats with an indirect draw call. ### Live Demo -[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +[![]()](https://christinaqiu3.com/Project5-WebGPU-Gaussian-Splat-Viewer/) ### Demo Video/GIF -[![](img/video.mp4)](TODO) +[![](images/hw_5_1.gif)](TODO) + +## Performance Analysis + +### Comparing point-cloud and gaussian renderer + +The point-cloud renderer displays each point as a simple dot or quad with uniform size and color, producing a sparse and noisy appearance without smooth transitions. It lacks depth-dependent opacity, shading, or blending between nearby points. + +The Gaussian renderer, produces smooth, continuous surfaces because each point is rendered as an anisotropic 2D Gaussian ellipse instead of a fixed-size point. The splats blend together, giving a more realistic and soft appearance, especially around regions with dense point samples. + +### Workgroup size affecting performance + +insert graphs + +### View-frustum culling give performance improvement -### (TODO: Your README) +insert graphs -*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. +### Number of guassians affects performance -This assignment has a considerable amount of performance analysis compared -to implementation work. Complete the implementation early to leave time! +insert graphs ### Credits diff --git a/Screenshot 2025-10-29 235258.png b/Screenshot 2025-10-29 235258.png new file mode 100644 index 0000000..6482b78 Binary files /dev/null and b/Screenshot 2025-10-29 235258.png differ diff --git a/images/Screenshot 2025-10-29 235015.png b/images/Screenshot 2025-10-29 235015.png new file mode 100644 index 0000000..aa22dd5 Binary files /dev/null and b/images/Screenshot 2025-10-29 235015.png differ diff --git a/images/Screenshot 2025-10-29 235258.png b/images/Screenshot 2025-10-29 235258.png new file mode 100644 index 0000000..6482b78 Binary files /dev/null and b/images/Screenshot 2025-10-29 235258.png differ diff --git a/images/hw_5_1.gif b/images/hw_5_1.gif new file mode 100644 index 0000000..21799d5 Binary files /dev/null and b/images/hw_5_1.gif differ diff --git a/src/renderers/gaussian-renderer.ts b/src/renderers/gaussian-renderer.ts index 1684523..501912d 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_settings_buffer: GPUBuffer, } // Utility to create GPU buffers @@ -17,7 +17,7 @@ const createBuffer = ( data?: ArrayBuffer | ArrayBufferView ) => { const buffer = device.createBuffer({ label, size, usage }); - if (data) device.queue.writeBuffer(buffer, 0, data); + if (data) device.queue.writeBuffer(buffer, 0, data as BufferSource); return buffer; }; @@ -36,6 +36,22 @@ export default function get_renderer( const nulling_data = new Uint32Array([0]); + const zero_buffer = createBuffer( + device, + 'zero buffer', + 4, + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + nulling_data + ); + + const render_settings_buffer = createBuffer( + device, + 'render settings buffer', + 8, + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + new Float32Array([1.0, pc.sh_deg]) + ); + // =============================================== // Create Compute Pipeline and Bind Groups // =============================================== @@ -63,24 +79,165 @@ export default function get_renderer( ], }); + const camera_bind_group = device.createBindGroup({ + label: 'camera bind group', + layout: preprocess_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer } }, + ], + }); + + const gaussian_preprocess_bind_group = device.createBindGroup({ + label: 'gaussian preprocess bind group', + layout: preprocess_pipeline.getBindGroupLayout(1), + entries: [ + { binding: 0, resource: { buffer: pc.gaussian_3d_buffer } }, + ], + }); + + var splat_size = 0; + splat_size += 16; // for position, size, color + splat_size *= pc.num_points; + + const splat_buffer = createBuffer( + device, + 'splat buffer', + splat_size, + GPUBufferUsage.STORAGE, + null + ); + + const compute_bind_group = device.createBindGroup({ + label: 'compute bind group', + layout: preprocess_pipeline.getBindGroupLayout(3), + entries: [ + { binding: 0, resource: { buffer: splat_buffer } },//buffer: pc.gaussian_3d_buffer } }, + { binding: 1, resource: { buffer: render_settings_buffer } }, + { binding: 2, resource: { buffer: pc.sh_buffer}}, + ], + }); // =============================================== // Create Render Pipeline and Bind Groups // =============================================== + // instance quads for each point in cloud, compute pipeline and then rendering pipeline + // use create buffer function tomake vertex and index buffer (using gpu buffer usage, buffer usage.indirect) + // pc.num_points number of quads + + const indirect_buffer = createBuffer( + device, + 'indirect draw buffer', + 4 * 4, // 4 uint32 values + GPUBufferUsage.INDIRECT | GPUBufferUsage.COPY_DST, + new Uint32Array([6, 0, 0, 0]) // 6 vertices per quad, num_points instances + ); + + const render_pipeline = device.createRenderPipeline({ + label: 'gaussian render pipeline', + layout: 'auto', + vertex: { + module: device.createShaderModule({ code: renderWGSL }), + entryPoint: 'vs_main', + buffers: [], + }, + fragment: { + module: device.createShaderModule({ code: renderWGSL }), + entryPoint: 'fs_main', + 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', + }, + } + }, + ], + }, + }); + const render_bind_group = device.createBindGroup({ + label: 'render bind group', + layout: render_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: splat_buffer } }, + { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer } }, + { binding: 2, resource: { buffer: camera_buffer}}, + ], + }); // =============================================== // Command Encoder Functions // =============================================== - + const compute = (encoder: GPUCommandEncoder) => { + const pass = encoder.beginComputePass(); + pass.setPipeline(preprocess_pipeline); + pass.setBindGroup(0, camera_bind_group); + pass.setBindGroup(1, gaussian_preprocess_bind_group); + pass.setBindGroup(2, sort_bind_group); + pass.setBindGroup(3, compute_bind_group); + pass.dispatchWorkgroups(Math.ceil(pc.num_points / C.histogram_wg_size)); + pass.end(); + } // =============================================== // Return Render Object // =============================================== return { frame: (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + + // clear sorter buffers + encoder.copyBufferToBuffer( + zero_buffer, + 0, + sorter.sort_info_buffer, + 0, + 4 + ); + encoder.copyBufferToBuffer( + zero_buffer, + 0, + sorter.sort_dispatch_indirect_buffer, + 0, + 4 + ); + + // Preprocess compute pass + compute(encoder); + sorter.sort(encoder); + + encoder.copyBufferToBuffer( + sorter.sort_info_buffer, + 0, + indirect_buffer, + 4, + 4 + ); + + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: texture_view, + loadOp: 'clear', + storeOp: 'store', + clearValue: { r: 0, g: 0, b: 0, a: 1 }, + }, + ], + }); + pass.setPipeline(render_pipeline); + pass.setBindGroup(0, render_bind_group); + pass.drawIndirect(indirect_buffer, 0); + pass.end(); + }, camera_buffer, + render_settings_buffer, }; } diff --git a/src/renderers/renderer.ts b/src/renderers/renderer.ts index ffdf9ba..71ba59b 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_settings_buffer, + 0, + new Float32Array([params.gaussian_multiplier]) //,0 + ); + } }); } diff --git a/src/shaders/gaussian.wgsl b/src/shaders/gaussian.wgsl index 759226d..58bc472 100644 --- a/src/shaders/gaussian.wgsl +++ b/src/shaders/gaussian.wgsl @@ -1,22 +1,102 @@ struct VertexOutput { @builtin(position) position: vec4, //TODO: information passed from vertex shader to fragment shader + @location(0) color: vec4, + @location(1) conic_opacity: vec4, + @location(2) center: vec2, }; struct Splat { //TODO: information defined in preprocess compute shader + pos: vec4, + // opacity: f32, + // rot: mat2x2, + scale: vec2, + color: array, + conic_opacity: array, }; +struct Camera { + viewMat: mat4x4, + invViewMat: mat4x4, + projMat: mat4x4, + invProjMat: mat4x4, + viewport: vec2, + focal: vec2, +} + +@group(0) @binding(0) +var splats : array; +@group(0) @binding(1) +var sort_indices : array; +@group(0) @binding(2) +var camera: Camera; + @vertex -fn vs_main( +fn vs_main( + @builtin(instance_index) global_instance_index: u32, + @builtin(vertex_index) local_vertex_index: u32 ) -> VertexOutput { - //TODO: reconstruct 2D quad based on information from splat, pass + //TODO: reconstruct 2D quad based on information from splat, pass to fragment shader + + let index = sort_indices[global_instance_index]; + let splat = splats[index]; + + let x = splat.pos.x; + let y = splat.pos.y; + let w = splat.scale.x * 2.0f; // because scale is half-width + let h = splat.scale.y * 2.0f; + +// array of 6 vec2s using x y w h +// does the order matter? + let quad = array,6>( + vec2(x - w, y + h), + vec2(x - w, y - h), + vec2(x + w, y - h), + vec2(x + w, y - h), + vec2(x + w, y + h), + vec2(x - w, y + h) + ); + +// vertex positions using local_vertex_index + let position = vec4(quad[local_vertex_index], 0.0, 1.0); // splat.pos.z for depth + + let unpacked_a = unpack2x16float(splat.color[0]); + let unpacked_b = unpack2x16float(splat.color[1]); + let color = vec4(unpacked_a.x, unpacked_a.y, unpacked_b.x, unpacked_b.y); + let size = (vec2f(w, h) * .5f + .5f) * camera.viewport.xy; + +// conic + let conic_unpacked_a = unpack2x16float(splat.conic_opacity[0]); + let conic_unpacked_b = unpack2x16float(splat.conic_opacity[1]); + let conic = vec3f(conic_unpacked_a.x, conic_unpacked_a.y, conic_unpacked_b.x); + let opacity = conic_unpacked_b.y; + var out: VertexOutput; - out.position = vec4(1. ,1. , 0., 1.); + out.position = position; + out.color = color; + out.conic_opacity = vec4f(conic, opacity); + out.center = vec2f(x, y); + return out; } @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4 { - return vec4(1.); + + var pos = (in.position.xy / camera.viewport) * 2.f - 1.f; + pos.y *= -1.f; + + var offset = (pos.xy - in.center.xy) * camera.viewport * vec2f(-.5f, .5f); + + var power = in.conic_opacity.x * pow(offset.x, 2.f) + in.conic_opacity.z * pow(offset.y, 2.f); + power = power * -.5f - in.conic_opacity.y * offset.x * offset.y; + + if (power > 0.f) { + return vec4f(0.f); + } + + let alpha = min(.99f, in.conic_opacity.w * exp(power)); + + 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..bbab815 100644 --- a/src/shaders/preprocess.wgsl +++ b/src/shaders/preprocess.wgsl @@ -57,9 +57,26 @@ struct Gaussian { struct Splat { //TODO: store information for 2D splat rendering + pos: vec4, + scale: vec2, + color: array, + conic_opacity: array, }; //TODO: bind your data here +// camera uniforms +@group(0) @binding(0) +var camera: CameraUniforms; +// render settings +// @group(0) @binding(1) +// var render_settings: RenderSettings; +// gaussians from load.ts +@group(1) @binding(0) +var gaussians : array; +// splats to be filled in preprocess compute shader +// @group(1) @binding(1) +// var splats : array; +// sorting related buffers @group(2) @binding(0) var sort_infos: SortInfos; @group(2) @binding(1) @@ -68,11 +85,30 @@ var sort_depths : array; var sort_indices : array; @group(2) @binding(3) var sort_dispatch: DispatchIndirect; +// storage for splats +@group(3) @binding(0) +var splats: array; +// render settings +@group(3) @binding(1) +var render_settings: RenderSettings; +@group(3) @binding(2) // 2 +var colors: array; /// 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) * 3 + c_idx % 2; + + if (c_idx%2 == 0) { + let unpacked_a = unpack2x16float(colors[index]); + let unpacked_b = unpack2x16float(colors[index + 1]); + return vec3f(unpacked_a.x, unpacked_a.y, unpacked_b.x); + } else { + let unpacked_a = unpack2x16float(colors[index]); + let unpacked_b = unpack2x16float(colors[index + 1]); + return vec3f(unpacked_a.y, unpacked_b.x, unpacked_b.y); + } } // spherical harmonics evaluation with Condon–Shortley phase @@ -111,8 +147,145 @@ 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 + // TODO: set up pipeline as described in instruction + + if (idx >= arrayLength(&gaussians)) { + return; + } + + let gaussian = gaussians[idx]; + let a = unpack2x16float(gaussian.pos_opacity[0]); + let b = unpack2x16float(gaussian.pos_opacity[1]); + let pos = vec4(a.x, a.y, b.x, 1.); + let opacity = b.y; + + // clip position + let clip_pos = camera.proj * camera.view * pos; + // + let ndc_pos = clip_pos.xyz / clip_pos.w; + + let view_depth = (camera.view * pos).z; + // frustum culling + if (ndc_pos.x < -1.2 || ndc_pos.x > 1.2 || ndc_pos.y < -1.2 || ndc_pos.y > 1.2 || view_depth < 0.0 || ndc_pos.z < 0.0 || ndc_pos.z > 1.2) { + return; + } + + let rot_a = unpack2x16float(gaussian.rot[0]); + let rot_b = unpack2x16float(gaussian.rot[1]); + let rot = vec4(rot_a.x, rot_a.y, rot_b.x, rot_b.y); + let rotMat = mat3x3( + vec3(1.0 - 2.0 * (rot.z * rot.z + rot.w * rot.w), 2.0 * (rot.y * rot.z - rot.x * rot.w), 2.0 * (rot.y * rot.w + rot.x * rot.z)), + vec3(2.0 * (rot.y * rot.z + rot.x * rot.w), 1.0 - 2.0 * (rot.y * rot.y + rot.w * rot.w), 2.0 * (rot.w * rot.z - rot.x * rot.y)), + vec3(2.0 * (rot.y * rot.w - rot.x * rot.z), 2.0 * (rot.z * rot.w + rot.x * rot.y), 1.0 - 2.0 * (rot.y * rot.y + rot.z * rot.z)) + ); + + let scale_a = unpack2x16float(gaussian.scale[0]); + let scale_b = unpack2x16float(gaussian.scale[1]); + let scale = exp(vec3(scale_a.x, scale_a.y, scale_b.x)); + + let scaleMat = mat3x3( + vec3(scale.x * render_settings.gaussian_scaling, 0.0, 0.0), + vec3(0.0, scale.y * render_settings.gaussian_scaling, 0.0), + vec3(0.0, 0.0, scale.z * render_settings.gaussian_scaling) + ); + +// covariance matrix + let covMat = transpose(scaleMat * rotMat) * (scaleMat * rotMat); + + let covariance = array( + covMat[0][0], + covMat[0][1], + covMat[0][2], + covMat[1][1], + covMat[1][2], + covMat[2][2] + ); + +// need to get t vector, j mat, w mat, t mat, v mat, to get 2d cov mat and cov + let temp = (camera.view * pos).xyz; + var t = temp; + +// .65 is a safe fraction of the visible field of view + let maxTx = max(.65f * camera.viewport.x / camera.focal.x, temp.x/temp.z); + let maxTy = max(.65f * camera.viewport.y / camera.focal.y, temp.y/temp.z); + t.x = min(.65f * camera.viewport.x / camera.focal.x, maxTx) * temp.z; + t.y = min(.65f * camera.viewport.y / camera.focal.y, maxTy) * temp.z; + +// let limx = 0.65 * camera.viewport.x / camera.focal.x; +// let limy = 0.65 * camera.viewport.y / camera.focal.y; +// let tx_over_z = temp.x / temp.z; +// let ty_over_z = temp.y / temp.z; +// t.x = min(limx, max(-limx, tx_over_z)) * temp.z; +// t.y = min(limy, max(-limy, ty_over_z)) * temp.z; + + + let jMat = mat3x3f( + camera.focal.x/t.z, 0.f, -(camera.focal.x * t.x) / (t.z * t.z), + 0.f, camera.focal.y/t.z, -(camera.focal.y * t.y) / (t.z * t.z), + 0.f, 0.f, 0.f + ); + + let wMat = transpose(mat3x3f(camera.view[0].xyz, camera.view[1].xyz, camera.view[2].xyz)); + + let tMat = wMat * jMat; + + let vMat = mat3x3f( + covariance[0], + covariance[1], + covariance[2], + covariance[1], + covariance[3], + covariance[4], + covariance[2], + covariance[4], + covariance[5], + + ); + + var covMat2D = transpose(tMat) * vMat * tMat; + covMat2D[0][0] += .3f; + covMat2D[1][1] += .3f; + // add so dont divide by 0 + + let covariance2D = vec3f(covMat2D[0][0], covMat2D[0][1], covMat2D[1][1]); + let keys_per_dispatch = workgroupSize * sortKeyPerThread; // increment DispatchIndirect.dispatchx each time you reach limit for one dispatch of keys + // TODO: compute splat size based on depth + var determinant = covariance2D.x * covariance2D.z - covariance2D.y * covariance2D.y; + if (determinant == 0.f) { return; } + + let middle = (covariance2D.x + covariance2D.z) * .5f; + let lambda1 = middle + sqrt(max(.1f, middle * middle - determinant)); + let lambda2 = middle - sqrt(max(.1f, middle * middle - determinant)); + let rad = ceil(3.f * sqrt(max(lambda1, lambda2))); + + let size = vec2f(rad, rad) / camera.viewport; + + let index = atomicAdd(&sort_infos.keys_size, 1u); + splats[index].pos = vec4(ndc_pos.x, ndc_pos.y, ndc_pos.z, 1.0); + splats[index].scale = size; + + + let direction = normalize(pos.xyz - camera.view_inv[3].xyz); + let color = computeColorFromSH(direction, idx, u32(render_settings.sh_deg)); + + splats[index].color[0] = pack2x16float(color.rg); + splats[index].color[1] = pack2x16float(vec2f(color.b, 1.f)); + + // conic + let conic = vec3f(covariance2D.z, -covariance2D.y, covariance2D.x) / determinant; + + splats[index].conic_opacity[0] = pack2x16float(conic.xy); + splats[index].conic_opacity[1] = pack2x16float(vec2f(conic.z, 1.f / (1.f + exp(-opacity)))); + + + + sort_depths[index] = bitcast(100.0f - view_depth); + sort_indices[index] = index; + + if (index % keys_per_dispatch == 0) { + atomicAdd(&sort_dispatch.dispatch_x, 1u); + } } \ No newline at end of file