diff --git a/README.md b/README.md index edffdaf..2f7da83 100644 --- a/README.md +++ b/README.md @@ -2,25 +2,44 @@ **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) +* Marcus Hedlund + * [LinkedIn](https://www.linkedin.com/in/marcushedlund/) +* Tested on: Windows 11, Intel Core Ultra 9 185H @ 2.5 GHz 16GB, NVIDIA GeForce RTX 4070 Laptop GPU 8GB (Personal Computer) + ### Live Demo -[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +[![](images/splat1.png)](http://mhedlund7.github.io/Project5-WebGPU-Gaussian-Splat-Viewer) ### Demo Video/GIF -[![](img/video.mp4)](TODO) +[![](images/bikegif.gif)](TODO) + +### Overview + +In this project I implement a real-time WebGPU Gaussian Splat Viewer that renders scenes using neural point clouds based on the paper [3D Gaussian Splatting for Real-Time Radiance Field Rendering](https://repo-sam.inria.fr/fungraph/3d-gaussian-splatting/). Scenes are loaded from .ply files and can be displayed as either point clouds or Gaussian splats. The Gaussian splat method uses 3D Gaussian distributions in order to represent geometry. It maps these 3D Gaussians to 2D screen-space conics using the Jacobian of the projection matrix, and performs View-Frustrum culling to only evaluate Gaussians that the camera can see. Additionally, during preprocessing it leverages spherical harmonics to evaluate color and sorts the splats by depth. Then in the rendering pass, we render a quad per splat and use the conic to apply exponential falloff for opacity and composite the Gaussians through alpha-blending to produce the final image. + +|![](images/pointcloud2.png)|![](images/splat2.png)| +|:--:|:--:| +| *Point Cloud Renderer* | *Gaussian Splat Renderer* | + +#### Comparing Point Cloud and Gaussian Renderer + +Visually, the point cloud method renders the scene as just that, a collection of points each of the same color. Surprisinngly this method is actually able to do a fairly good job of encoding what is represented in the scene as can be seen in the above image, but it still leaves much to be desired. Gaussian splatting on the other hand is able to leverage the process to fully recreate a scene. It has accurate color, texture, and overall has an extrextremelymely realistic appearance. Performance-wise though as expected the point cloud method performs much better. In this case for the above scene I got 4.11 ms/frame for the point cloud and 6.58 ms/frame for the Gaussian Splasplattingtting. This makes a lot of sense because of how much more involved the Gaussian splatting is, having to perform many precomputations like 3D covariance, 2D conics, color, sorting, and opacity falloff instead of just creating points. + +#### Workgroup Size Effect on Performance + +Decreasing the workgroup size will generally lower the performance of Gaussian splatting. This is because it makes it so not as many splats can be processed in parallel, slowing down the computation. For me, I achieved maximum performance at a workgroup of size 256. This is likely because at higher workgroup sizes there is too high of an overhead in both dispatch time and memory pressure, and the GPU will already be saturated so the additional workgroups don't fully leverage additional parallelism, resulting in slower computation. + +#### View-Frustrum Culling Effect on Performance + +I additionally performed view-frustrum culling during preprocessing by creating a bounding box around the camera's frustrum (extended by 1.2x in x and y to account for gaussians slightly outside the frustrum that could still influence the image) and not processing gaussians that fall outside that range. This was able to greatly improve performance when the camera was positioned closer to the scene so that not all of the gaussians were in sight, however there was close to no effect when the camera was positioned farther away so that all gaussians were visible. This makes sense because when we are actually able to cull Gaussians and not process them through View-Frustrum Culling it just leads to us having to process fewer items in each of the subsequent steps, leading to quicker computation time. + +#### Number of Gaussians Effect on Performance -### (TODO: Your README) +We are able to compare the effect the number of Gaussians has on performance by looking at different scenes with a higher or lower amount of points in them. As can be expected increasing the number of Gaussians in a scene decreases the performance of our viewer. This is because we will have a larger number of Gaussians to compute information for and process, leading to higher compute time and lower 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. -This assignment has a considerable amount of performance analysis compared -to implementation work. Complete the implementation early to leave time! ### Credits diff --git a/images/bikegif.gif b/images/bikegif.gif new file mode 100644 index 0000000..ea3cba4 Binary files /dev/null and b/images/bikegif.gif differ diff --git a/images/pointcloud1.png b/images/pointcloud1.png new file mode 100644 index 0000000..c31dfac Binary files /dev/null and b/images/pointcloud1.png differ diff --git a/images/pointcloud2.png b/images/pointcloud2.png new file mode 100644 index 0000000..b2abf01 Binary files /dev/null and b/images/pointcloud2.png differ diff --git a/images/splat1.png b/images/splat1.png new file mode 100644 index 0000000..1235a1e Binary files /dev/null and b/images/splat1.png differ diff --git a/images/splat2.png b/images/splat2.png new file mode 100644 index 0000000..80e2778 Binary files /dev/null and b/images/splat2.png differ diff --git a/package-lock.json b/package-lock.json index 04843bd..694c409 100644 --- a/package-lock.json +++ b/package-lock.json @@ -12,6 +12,7 @@ "@loaders.gl/ply": "^4.2.2", "@petamoriken/float16": "^3.8.7", "tweakpane": "^3.1.8", + "tweakpane-plugin-file-import": "^0.2.0", "wgpu-matrix": "^3.2.0" }, "devDependencies": { diff --git a/src/renderers/gaussian-renderer.ts b/src/renderers/gaussian-renderer.ts index 1684523..0c2ac87 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 { - + setScalingMultiplier: (multiplier: number) => void; } // Utility to create GPU buffers @@ -36,6 +36,47 @@ export default function get_renderer( const nulling_data = new Uint32Array([0]); + // settings for gaussian rendering (scaling, sh_deg) + const gaussian_rendering_settings_array = new Float32Array(4); + gaussian_rendering_settings_array[0] = 1.0; // scaling multiplier + gaussian_rendering_settings_array[1] = pc.sh_deg; // padding + + const gaussian_rendering_settings_buffer = createBuffer( + device, + 'gaussian uniforms', + gaussian_rendering_settings_array.byteLength, + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + gaussian_rendering_settings_array + ); + + // splat buffer + const splatSize = 56; + const splatBuffer = createBuffer( + device, + 'splat buffer', + pc.num_points * splatSize, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + ); + + // indirect draw + const drawArgsArray = new Uint32Array([6, 0, 0, 0]); // vertexCount, instanceCount, firstVertex, firstInstance + const indirectDrawBuffer = createBuffer( + device, + 'indirect draw buffer', + drawArgsArray.byteLength, + GPUBufferUsage.INDIRECT | GPUBufferUsage.COPY_DST, + drawArgsArray, + ); + + // nulling buffer + const nulling_buffer = createBuffer( + device, + 'nulling buffer', + nulling_data.byteLength, + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + nulling_data, + ); + // =============================================== // Create Compute Pipeline and Bind Groups // =============================================== @@ -52,6 +93,26 @@ export default function get_renderer( }, }); + + const preprocess_camera_bind_group = device.createBindGroup({ + label: 'camera bind group', + layout: preprocess_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer } }, + ], + }); + + const preprocess_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: splatBuffer } }, + { binding: 2, resource: { buffer: gaussian_rendering_settings_buffer } }, + { binding: 3, resource: { buffer: pc.sh_buffer } } + ], + }); + const sort_bind_group = device.createBindGroup({ label: 'sort', layout: preprocess_pipeline.getBindGroupLayout(2), @@ -67,20 +128,132 @@ export default function get_renderer( // =============================================== // Create Render Pipeline and Bind Groups // =============================================== - + const render_pipeline = device.createRenderPipeline({ + label: 'gaussian render', + layout: 'auto', + vertex: { + module: device.createShaderModule( + { + code: renderWGSL + } + ), + entryPoint: 'vs_main', + }, + 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_gaussian_bind_group = device.createBindGroup({ + label: 'render gaussian bind group', + layout: render_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: splatBuffer } }, + { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer } }, + { binding: 2, resource: { buffer: camera_buffer } } + ], + }); // =============================================== // Command Encoder Functions // =============================================== + const preprocess = (encoder: GPUCommandEncoder) => { + const pass = encoder.beginComputePass({ + label: 'preprocess pass', + }); + pass.setPipeline(preprocess_pipeline); + pass.setBindGroup(0, preprocess_camera_bind_group); + pass.setBindGroup(1, preprocess_gaussian_bind_group); + pass.setBindGroup(2, sort_bind_group); + const numWorkgroups = Math.ceil(pc.num_points / C.histogram_wg_size); + + pass.dispatchWorkgroups(numWorkgroups); + pass.end(); + }; + + + const render = (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + const pass = encoder.beginRenderPass({ + label: 'gaussian render pass', + 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_gaussian_bind_group); + pass.drawIndirect(indirectDrawBuffer, 0); + pass.end(); + }; + + // =============================================== + // Scaling Multiplier Function + // =============================================== + + // update scaling multiplier + function setScalingMultiplier(multiplier: number) { + gaussian_rendering_settings_array[0] = multiplier; + device.queue.writeBuffer( + gaussian_rendering_settings_buffer, 0, + gaussian_rendering_settings_array + ); + } // =============================================== // 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, + indirectDrawBuffer, 4, 4 + ); + + render(encoder, texture_view); }, camera_buffer, + setScalingMultiplier, }; } diff --git a/src/renderers/renderer.ts b/src/renderers/renderer.ts index ffdf9ba..61878d7 100644 --- a/src/renderers/renderer.ts +++ b/src/renderers/renderer.ts @@ -122,6 +122,9 @@ export default async function init( {min: 0, max: 1.5} ).on('change', (e) => { //TODO: Bind constants to the gaussian renderer. + if (gaussian_renderer) { + gaussian_renderer.setScalingMultiplier(e.value); + } }); } diff --git a/src/shaders/gaussian.wgsl b/src/shaders/gaussian.wgsl index 759226d..86c0762 100644 --- a/src/shaders/gaussian.wgsl +++ b/src/shaders/gaussian.wgsl @@ -1,22 +1,86 @@ struct VertexOutput { @builtin(position) position: vec4, //TODO: information passed from vertex shader to fragment shader + @location(0) color : vec4, + @location(1) conic : vec3, + @location(2) opacity : f32, + @location(3) center : vec2 }; struct Splat { - //TODO: information defined in preprocess compute shader + //TODO: store information for 2D splat rendering + ndc_center: vec2, + radius: vec2, + ndc_depth: f32, + pad: f32, + color: vec4, + conic_opacity: vec4, }; +struct CameraUniforms { + view: mat4x4, + view_inv: mat4x4, + proj: mat4x4, + proj_inv: mat4x4, + viewport: vec2, + focal: vec2 +}; + +@group(0) @binding(0) +var splats: array; +@group(0) @binding(1) +var sorted_indices : array; +@group(0) @binding(2) +var camera: CameraUniforms; + +fn quad_vertex_offset(vertex_index: u32) -> vec2 { + var offsets = 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 offsets[vertex_index]; +} + @vertex -fn vs_main( +fn vs_main(@builtin(vertex_index) vertex_index: u32, @builtin(instance_index) instance_index: 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 = sorted_indices[instance_index]; + + // splat info + let splat = splats[splat_idx]; + let offset = quad_vertex_offset(vertex_index % 6u); + let ndc_offset = vec2( + offset.x * splat.radius.x, + offset.y * splat.radius.y + ); + let ndc_out = splat.ndc_center + ndc_offset; + + // out info + out.position = vec4(ndc_out.x, ndc_out.y, splat.ndc_depth, 1.0); + out.color = splat.color; + out.conic = splat.conic_opacity.xyz; + out.opacity = splat.conic_opacity.w; + // convert center into pixel space + out.center = (0.5 + splat.ndc_center * vec2(0.5, -0.5)) * camera.viewport; return out; } @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4 { - return vec4(1.); + var dist = in.position.xy - in.center.xy; + dist.x = -dist.x; + let power = -0.5 * (in.conic.x * dist.x * dist.x + in.conic.z * dist.y * dist.y) - in.conic.y * dist.x * dist.y; + if (power > 0.0) { + return vec4(0.0, 0.0, 0.0, 0.0); + } + + let alpha = min(0.99, in.opacity * 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..4b38fba 100644 --- a/src/shaders/point_cloud.wgsl +++ b/src/shaders/point_cloud.wgsl @@ -35,7 +35,9 @@ fn vs_main( let pos = vec4(a.x, a.y, b.x, 1.); // TODO: MVP calculations - out.position = pos; + let view_pos = camera.view * pos; + let clip_pos = camera.proj * view_pos; + out.position = clip_pos; return out; } diff --git a/src/shaders/preprocess.wgsl b/src/shaders/preprocess.wgsl index bbc63f5..91d921f 100644 --- a/src/shaders/preprocess.wgsl +++ b/src/shaders/preprocess.wgsl @@ -57,9 +57,27 @@ struct Gaussian { struct Splat { //TODO: store information for 2D splat rendering + ndc_center: vec2, + radius: vec2, + ndc_depth: f32, + pad: f32, + color: vec4, + conic_opacity: vec4, }; //TODO: bind your data here + +@group(0) @binding(0) +var camera: CameraUniforms; +@group(1) @binding(0) +var gaussians : array; +@group(1) @binding(1) +var splats: array; +@group(1) @binding(2) +var render_settings: RenderSettings; +@group(1) @binding(3) +var sh_buffer: array; + @group(2) @binding(0) var sort_infos: SortInfos; @group(2) @binding(1) @@ -69,12 +87,121 @@ var sort_indices : array; @group(2) @binding(3) var sort_dispatch: DispatchIndirect; +// see if ndc point is inside bounding box for 1.2x frustum +fn in_ndc_bounding_box(ndc_pos: vec3) -> bool { + let s = 1.2; + return ndc_pos.x >= -s && ndc_pos.x <= s && + ndc_pos.y >= -s && ndc_pos.y <= s && + ndc_pos.z > -0.0; +} + /// 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 idx = splat_idx * 24 + 3 * c_idx / 2 + c_idx % 2; + + let colrg = unpack2x16float(sh_buffer[idx]); + let colba = unpack2x16float(sh_buffer[idx + 1]); + + if (c_idx % 2 == 0) { + return vec3(colrg.x, colrg.y, colba.x); + } else { + return vec3(colrg.y, colba.x, colba.y); + } +} + +fn mat3_from_quat(q: vec4) -> mat3x3 { + let x = q.x; + let y = q.y; + let z = q.z; + let r = q.w; + return mat3x3( + vec3(1.0 - 2.0 * (y * y + z * z), 2.0 * (x * y - z * r), 2.0 * (x * z + y * r)), + vec3(2.0 * (x * y + z * r), 1.0 - 2.0 * (x * x + z * z), 2.0 * (y * z - x * r)), + vec3(2.0 * (x * z - y * r), 2.0 * (y * z + x * r), 1.0 - 2.0 * (x * x + y * y)) + ); +} + +fn compute3DCovarianceFromGaussian(g: Gaussian, modifier: f32) -> mat3x3 { + let scale_xy = unpack2x16float(g.scale[0]); + let scale_zw = unpack2x16float(g.scale[1]); + let scales = vec3( + exp(scale_xy.x), + exp(scale_xy.y), + exp(scale_zw.x) + ); + + // Debug: clamp scales to reasonable values + let rot_xy = unpack2x16float(g.rot[0]); + let rot_zr = unpack2x16float(g.rot[1]); + let quat = vec4(rot_xy.x, rot_xy.y, rot_zr.x, rot_zr.y); + let rot_mat = mat3_from_quat(quat); + + let scale_mat = mat3x3( + vec3(modifier * scales.x, 0.0, 0.0), + vec3(0.0, modifier * scales.y, 0.0), + vec3(0.0, 0.0, modifier * scales.z) + ); + + let covariance = rot_mat * scale_mat * transpose(scale_mat) * transpose(rot_mat); + return covariance; } +fn compute2DCovarianceFromGaussian(g: Gaussian, view_pos3: vec3, modifier: f32, cam: CameraUniforms) -> mat3x3 { + let cov3D = compute3DCovarianceFromGaussian(g, modifier); + + + let fx = cam.focal.x; + let fy = cam.focal.y; + + let jacobian = mat3x3( + vec3(fx / view_pos3.z, 0.0, -fx * view_pos3.x / (view_pos3.z * view_pos3.z)), + vec3(0.0, fy / view_pos3.z, -fy * view_pos3.y / (view_pos3.z * view_pos3.z)), + vec3(0.0, 0.0, 0.0) + ); + + let W = transpose(mat3x3(cam.view[0].xyz, cam.view[1].xyz, cam.view[2].xyz)); + + let T = W * jacobian; + + let Vrk = 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(T) * Vrk * T; + + cov2D[0][0] += 0.3; + cov2D[1][1] += 0.3; + + return cov2D; +} + +fn computeRadiusFromCovariance(cov2D: mat3x3) -> f32 { + let det = cov2D[0][0] * cov2D[1][1] - cov2D[0][1] * cov2D[1][0]; + let mid = 0.5 * (cov2D[0][0] + cov2D[1][1]); + let lambdaOffset = sqrt(max(0.1, mid * mid - det)); + + let lambda1 = mid + lambdaOffset; + let lambda2 = mid - lambdaOffset; + + let max_lambda = max(lambda1, lambda2); + + let radius = ceil(3.0 * sqrt(max_lambda)); + return radius; +} + +fn computeConicFromCovariance(cov2D: mat3x3) -> mat2x2 { + let det = cov2D[0][0] * cov2D[1][1] - cov2D[0][1] * cov2D[1][0]; + let det_inv = 1.0 / det; + return det_inv * mat2x2( + cov2D[1][1], -cov2D[0][1], + -cov2D[1][0], cov2D[0][0] + ); +} + + // spherical harmonics evaluation with Condon–Shortley phase fn computeColorFromSH(dir: vec3, v_idx: u32, sh_deg: u32) -> vec3 { var result = SH_C0 * sh_coef(v_idx, 0u); @@ -105,7 +232,7 @@ fn computeColorFromSH(dir: vec3, v_idx: u32, sh_deg: u32) -> vec3 { } result += 0.5; - return max(vec3(0.), result); + return max(vec3(0.), result); } @compute @workgroup_size(workgroupSize,1,1) @@ -113,6 +240,67 @@ fn preprocess(@builtin(global_invocation_id) gid: vec3, @builtin(num_workgr let idx = gid.x; //TODO: set up pipeline as described in instruction + let n = arrayLength(&gaussians); + if (idx >= n) { + return; + } + + let g = gaussians[idx]; + // stored as x, y, z, opacity + let a = unpack2x16float(g.pos_opacity[0]); + let b = unpack2x16float(g.pos_opacity[1]); + let world_pos = vec4(a.x, a.y, b.x, 1.0); + let opacity = 1.0 / (1.0 + exp(-b.y)); + + // convert from world to ndc space + let view_pos = camera.view * world_pos; + let clip_pos = camera.proj * view_pos; + let ndc_pos = clip_pos.xyz / clip_pos.w; + + // discard splats outside of ndc bounding box + if (!in_ndc_bounding_box(ndc_pos)) { + return; + } + + // fill in splat info + var splat: Splat; + splat.ndc_center = ndc_pos.xy; + + // use quad to represent splat for now + // evaluate radius + + let view_pos3 = (camera.view * world_pos).xyz; + let cov2D = compute2DCovarianceFromGaussian(g, view_pos3, render_settings.gaussian_scaling, camera); + let conic = computeConicFromCovariance(cov2D); + let conic_opacity = vec4(conic[0][0], conic[0][1], conic[1][1], opacity); + let radius = computeRadiusFromCovariance(cov2D); + + + let radius2D = radius * vec2(1.0, 1.0) / camera.viewport; + splat.radius = radius2D; // assuming square viewport + splat.ndc_depth = ndc_pos.z; + splat.conic_opacity = conic_opacity; + + // evaluate color + + + let sh_deg = u32(render_settings.sh_deg); + let camera_pos = -camera.view[3].xyz; + + let view_dir = normalize(world_pos.xyz - camera_pos); + let color = computeColorFromSH(view_dir, idx, sh_deg); + splat.color = vec4(color, 1.0); + + let splat_idx = atomicAdd(&sort_infos.keys_size, 1); + + splats[splat_idx] = splat; + + sort_depths[splat_idx] = bitcast(100.0 - view_pos.z); + sort_indices[splat_idx] = splat_idx; + let keys_per_dispatch = workgroupSize * sortKeyPerThread; // increment DispatchIndirect.dispatchx each time you reach limit for one dispatch of keys + if (splat_idx % keys_per_dispatch == 0) { + atomicAdd(&sort_dispatch.dispatch_x, 1); + } } \ No newline at end of file