diff --git a/.gitignore b/.gitignore index 054e565..1fcae35 100644 --- a/.gitignore +++ b/.gitignore @@ -24,4 +24,6 @@ dist-ssr *.sw? /.vite -*/scenes \ No newline at end of file +*/scenes +scene/ +*/package-lock.json \ No newline at end of file diff --git a/README.md b/README.md index f99cdff..a28f0b3 100644 --- a/README.md +++ b/README.md @@ -1,31 +1,37 @@ -# Project5-WebGPU-Gaussian-Splat-Viewer +**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 5 - WebGPU Gaussian Splat Viewer** -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 4** +![main](images/main.png) -* (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) +* Yin Tang + * [Linkedin](https://www.linkedin.com/in/yin-tang-jackeyty/), [Github](https://github.com/JackeyTY), [Personal Website](https://jackeytang.com/) +* Tested on: Google Chrome 130.0 on Windows 11 Pro, AMD Ryzen 9 7950X @ 5.00GHz 64GB, NVIDIA GeForce RTX 4090 24GB (personal desktop) -### Live Demo +
-[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +### Overview -### Demo Video/GIF +In this project, I implemented a **3D Gaussian Splat Viewer** using **WebGPU**, which is based on the paper [3D Gaussian Splatting for Real-Time Radiance Field Rendering](https://repo-sam.inria.fr/fungraph/3d-gaussian-splatting/). Online interactive [demo](https://www.jackeytang.com/WebGPU-Gaussian-Splat-Viewer/), remember to download scene files first. -[![](img/video.mp4)](TODO) +**Gaussian Splatting** is a novel technique in computer graphics and neural rendering that represents a scene or object using a collection of 3D Gaussian functions instead of traditional polygonal meshes or voxel grids. Each Gaussian, or "splat," has a position, size, orientation, and opacity, defining a smooth, continuous volume that approximates the underlying surface or appearance. This approach allows for efficient rendering, as the splats can be rasterized quickly, and their smooth nature helps reduce aliasing artifacts. Gaussian splatting is particularly well-suited for applications in real-time rendering, neural radiance fields (NeRFs), and point-based rendering, where it offers a compact and differentiable representation. By leveraging GPU acceleration, it enables high-quality rendering with fewer computational resources, making it a promising alternative for interactive graphics, virtual reality, and dynamic scene representations. -### (TODO: Your README) +![next](images/next.png) -*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! +### Performance Analysis -### Credits +- **Compare your results from point-cloud and gaussian renderer, what are the differences?** -- [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) +​ Point-cloud is much easier to draw since only individual points are rendered, whereas gaussian splats render consists of more computation in the computer and render passes, which takes comparably large amount of time. On the other hand, the quality of the image rendered by the gaussian render is significantly better than that of point-cloud. + +- **For gaussian renderer, how does changing the workgroup-size affect performance? Why do you think this is?** + +​ Since the maximum total number of invocations per workgroup is 256, we can only decrease the workgroup size in the x-dimension, which results is worse performance since more workgroups need to be dispatched and less utilization. + +- **Does view-frustum culling give performance improvement? Why do you think this is?** + +​ It does improve the performance by around 5%, which is because we don't need to process out-of-bound splats so less computation and draw calls, but given the size of the scene, it is not evident. + +- **Does number of guassians affect performance? Why do you think this is?** + +​ As the number of gaussians, which is the upper bound for draw invocations, the performance decreases as more computation and draw calls need to be processed. diff --git a/images/main.png b/images/main.png new file mode 100644 index 0000000..08f722f Binary files /dev/null and b/images/main.png differ diff --git a/images/next.png b/images/next.png new file mode 100644 index 0000000..7a67d86 Binary files /dev/null and b/images/next.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..30d1e6a 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 { - + scaling_buffer: GPUBuffer } // Utility to create GPU buffers @@ -34,13 +34,48 @@ export default function get_renderer( // Initialize GPU Buffers // =============================================== + const scaling_buffer = createBuffer( + device, + 'render_settings_buffer', + 4, + GPUBufferUsage.COPY_DST | GPUBufferUsage.UNIFORM, + new Float32Array([1.0]) + ); + const nulling_data = new Uint32Array([0]); + const nulling_buffer = createBuffer( + device, + 'nulling buffer', + 4, + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + nulling_data + ); + + const splat_size = 24; + + const splat_buffer = createBuffer( + device, + 'splat buffer', + splat_size * pc.num_points, + GPUBufferUsage.STORAGE, + null + ); + + const indirect_buffer = createBuffer( + device, + 'indirect buffer', + 20, + GPUBufferUsage.COPY_DST | GPUBufferUsage.INDIRECT, + new Uint32Array([6, 0, 0, 0, 0]) + ); + // =============================================== // Create Compute Pipeline and Bind Groups // =============================================== + const preprocess_pipeline = device.createComputePipeline({ - label: 'preprocess', + label: 'preprocess pipeline', layout: 'auto', compute: { module: device.createShaderModule({ code: preprocessWGSL }), @@ -48,39 +83,169 @@ export default function get_renderer( constants: { workgroupSize: C.histogram_wg_size, sortKeyPerThread: c_histogram_block_rows, - }, - }, + shDegree: pc.sh_deg + } + } + }); + + const camera_bind_group = device.createBindGroup({ + label: 'preprocess camera bind group', + layout: preprocess_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer } } + ] + }); + + const gaussian_bind_group = device.createBindGroup({ + label: 'preprocess gaussians bind group', + layout: preprocess_pipeline.getBindGroupLayout(1), + entries: [ + { binding: 0, resource: { buffer: pc.gaussian_3d_buffer } } + ] }); const sort_bind_group = device.createBindGroup({ - label: 'sort', + label: 'preprocess sort bind group', layout: preprocess_pipeline.getBindGroupLayout(2), entries: [ { binding: 0, resource: { buffer: sorter.sort_info_buffer } }, { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_depths_buffer } }, { binding: 2, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer } }, - { binding: 3, resource: { buffer: sorter.sort_dispatch_indirect_buffer } }, - ], + { binding: 3, resource: { buffer: sorter.sort_dispatch_indirect_buffer } } + ] }); + const splat_preprocess_bind_group = device.createBindGroup({ + label: 'preprocess splat bind group', + layout: preprocess_pipeline.getBindGroupLayout(3), + entries: [ + { binding: 0, resource: { buffer: splat_buffer } }, + { binding: 1, resource: { buffer: scaling_buffer } }, + { binding: 2, resource: { buffer: pc.sh_buffer } } + ] + }); // =============================================== // Create Render Pipeline and Bind Groups // =============================================== - + + const gaussian_render_pipeline = device.createRenderPipeline({ + label: 'gaussian render pipeline', + 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 splat_render_bind_group = device.createBindGroup({ + label: 'gaussian render splat bind group', + layout: gaussian_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 preprocess_pass = (encoder: GPUCommandEncoder) => { + const preprocess_pass = encoder.beginComputePass(); + + preprocess_pass.setPipeline(preprocess_pipeline); + + preprocess_pass.setBindGroup(0, camera_bind_group); + preprocess_pass.setBindGroup(1, gaussian_bind_group); + preprocess_pass.setBindGroup(2, sort_bind_group); + preprocess_pass.setBindGroup(3, splat_preprocess_bind_group); + + preprocess_pass.dispatchWorkgroups(Math.ceil(pc.num_points / C.histogram_wg_size)); + + preprocess_pass.end(); + }; + const gaussian_render_pass = (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + const gaussian_pass = encoder.beginRenderPass({ + label: 'gaussian render', + colorAttachments: [ + { + view: texture_view, + loadOp: 'clear', + storeOp: 'store', + clearValue: [0, 0, 0, 1] + } + ], + }); + + gaussian_pass.setPipeline(gaussian_render_pipeline); + gaussian_pass.setBindGroup(0, splat_render_bind_group); + + gaussian_pass.drawIndirect(indirect_buffer, 0); + gaussian_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_pass(encoder); + sorter.sort(encoder); + + encoder.copyBufferToBuffer( + sorter.sort_info_buffer, + 0, + indirect_buffer, + 4, + 4 + ); + + gaussian_render_pass(encoder, texture_view); + }, + camera_buffer, + scaling_buffer }; } diff --git a/src/renderers/renderer.ts b/src/renderers/renderer.ts index ffdf9ba..f2fed26 100644 --- a/src/renderers/renderer.ts +++ b/src/renderers/renderer.ts @@ -121,7 +121,9 @@ export default async function init( 'gaussian_multiplier', {min: 0, max: 1.5} ).on('change', (e) => { - //TODO: Bind constants to the gaussian renderer. + if (gaussian_renderer) { + device.queue.writeBuffer(gaussian_renderer.scaling_buffer, 0, new Float32Array([params.gaussian_multiplier])); + } }); } diff --git a/src/shaders/gaussian.wgsl b/src/shaders/gaussian.wgsl index 759226d..b1d03b2 100644 --- a/src/shaders/gaussian.wgsl +++ b/src/shaders/gaussian.wgsl @@ -1,22 +1,93 @@ struct VertexOutput { @builtin(position) position: vec4, - //TODO: information passed from vertex shader to fragment shader + @location(0) wh: vec2f, + @location(1) color: vec4f, + @location(2) conic_opa: vec4f, + @location(3) center: vec2f +}; + +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 + xy: u32, + wh: u32, + rg: u32, + ba: u32, + co: u32, + cp: u32 }; +@group(0) @binding(0) +var splats: array; +@group(0) @binding(1) +var sort_indices : array; +@group(0) @binding(2) +var camera: CameraUniforms; + @vertex fn vs_main( + @builtin(instance_index) instanceIndex: u32, + @builtin(vertex_index) vertexIndex: u32 ) -> VertexOutput { - //TODO: reconstruct 2D quad based on information from splat, pass var out: VertexOutput; - out.position = vec4(1. ,1. , 0., 1.); + + let index = sort_indices[instanceIndex]; + let splat = splats[index]; + let xy = unpack2x16float(splat.xy); + let wh = unpack2x16float(splat.wh); + + let corners = array( + vec2f(xy.x - wh.x, xy.y + wh.y), + vec2f(xy.x - wh.x, xy.y - wh.y), + vec2f(xy.x + wh.x, xy.y - wh.y), + vec2f(xy.x + wh.x, xy.y - wh.y), + vec2f(xy.x + wh.x, xy.y + wh.y), + vec2f(xy.x - wh.x, xy.y + wh.y), + ); + let pos = vec4(corners[vertexIndex].x, corners[vertexIndex].y, 0, 1); + + let rg = unpack2x16float(splat.rg); + let ba = unpack2x16float(splat.ba); + let color = vec4f(rg.x, rg.y, ba.x, ba.y); + + let co = unpack2x16float(splat.co); + let cp = unpack2x16float(splat.cp); + let conic_opa = vec4f(co.x, co.y, cp.x, cp.y); + + out.position = pos; + out.wh = wh; + out.color = color; + out.conic_opa = conic_opa; + out.center = xy; + return out; } @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4 { - return vec4(1.); + var xy = (in.position.xy / camera.viewport) * 2.0f - 1.0f; + xy.y *= -1.0f; + + var d = xy - in.center; + d.x *= -1.0f; + d *= camera.viewport * 0.5f; + + let power = -0.5f * (in.conic_opa.x * d.x * d.x + in.conic_opa.z * d.y * d.y) - in.conic_opa.y * d.x * d.y; + + if (power > 0.0f) { + return vec4f(0.0f, 0.0f, 0.0f, 0.0f); + } + + let alpha = min(0.99f, in.conic_opa.w * exp(power)); + + return in.color * alpha; + //return vec4(in.wh.x, in.wh.y, 0.0, 1.0); + //return in.color; } \ No newline at end of file diff --git a/src/shaders/point_cloud.wgsl b/src/shaders/point_cloud.wgsl index 01dded1..ee5c067 100644 --- a/src/shaders/point_cloud.wgsl +++ b/src/shaders/point_cloud.wgsl @@ -34,8 +34,7 @@ fn vs_main( let b = unpack2x16float(vertex.pos_opacity[1]); 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..adc440b 100644 --- a/src/shaders/preprocess.wgsl +++ b/src/shaders/preprocess.wgsl @@ -19,6 +19,7 @@ const SH_C3 = array( override workgroupSize: u32; override sortKeyPerThread: u32; +override shDegree: u32; struct DispatchIndirect { dispatch_x: atomic, @@ -44,11 +45,6 @@ struct CameraUniforms { focal: vec2 }; -struct RenderSettings { - gaussian_scaling: f32, - sh_deg: f32, -} - struct Gaussian { pos_opacity: array, rot: array, @@ -56,10 +52,20 @@ struct Gaussian { }; struct Splat { - //TODO: store information for 2D splat rendering + xy: u32, + wh: u32, + rg: u32, + ba: u32, + co: u32, + cp: u32 }; -//TODO: bind your data here +@group(0) @binding(0) +var camera: CameraUniforms; + +@group(1) @binding(0) +var gaussians: array; + @group(2) @binding(0) var sort_infos: SortInfos; @group(2) @binding(1) @@ -69,13 +75,27 @@ var sort_indices : array; @group(2) @binding(3) var sort_dispatch: DispatchIndirect; -/// reads the ith sh coef from the storage buffer +@group(3) @binding(0) +var splats: array; +@group(3) @binding(1) +var scaling: f32; +@group(3) @binding(2) +var shs: array; + 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 i = splat_idx * 24 + (c_idx / 2) * 3 + c_idx % 2; + + if (c_idx % 2 == 0) { + let rg = unpack2x16float(shs[i + 0]); + let br = unpack2x16float(shs[i + 1]); + return vec3f(rg.x, rg.y, br.x); + } else { + let br = unpack2x16float(shs[i + 0]); + let gb = unpack2x16float(shs[i + 1]); + return vec3f(br.y, gb.x, gb.y); + } } -// 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); @@ -108,11 +128,152 @@ fn computeColorFromSH(dir: vec3, v_idx: u32, sh_deg: u32) -> vec3 { return max(vec3(0.), result); } +fn quatToRotationMatrix(rot: vec4f) -> mat3x3f { + let r = rot.x; + let x = rot.y; + let y = rot.z; + let z = rot.w; + + return mat3x3f( + vec3f(1.0 - 2.0 * (y * y + z * z), 2.0 * (x * y - r * z), 2.0 * (x * z + r * y)), + vec3f(2.0 * (x * y + r * z), 1.0 - 2.0 * (x * x + z * z), 2.0 * (y * z - r * x)), + vec3f(2.0 * (x * z - r * y), 2.0 * (y * z + r * x), 1.0 - 2.0 * (x * x + y * y)) + ); +} + +fn computeCov3D(scale: vec3f, scaling: f32, rot: vec4f) -> array { + let S = mat3x3f( + vec3f(scaling * scale.x, 0.0, 0.0), + vec3f(0.0, scaling * scale.y, 0.0), + vec3f(0.0, 0.0, scaling * scale.z) + ); + + let R = quatToRotationMatrix(rot); + + let M = S * R; + + let Sigma = transpose(M) * M; + + var cov3D: array; + cov3D[0] = Sigma[0][0]; + cov3D[1] = Sigma[0][1]; + cov3D[2] = Sigma[0][2]; + cov3D[3] = Sigma[1][1]; + cov3D[4] = Sigma[1][2]; + cov3D[5] = Sigma[2][2]; + + return cov3D; +} + +fn computeCov2D(pos: vec4f, focal_x: f32, focal_y: f32, tan_fovx: f32, tan_fovy: f32, cov3D: array, view: mat4x4f) -> vec3f { + var t = (view * pos).xyz; + + let limx = 1.3f * tan_fovx; + let limy = 1.3f * tan_fovy; + let txtz = t.x / t.z; + let tytz = t.y / t.z; + t.x = min(limx, max(-limx, txtz)) * t.z; + t.y = min(limy, max(-limy, tytz)) * t.z; + + let J = mat3x3f( + vec3f(focal_x / t.z, 0.0, -(focal_x * t.x) / (t.z * t.z)), + vec3f(0.0, focal_y / t.z, -(focal_y * t.y) / (t.z * t.z)), + vec3f(0.0, 0.0, 0.0) + ); + + let W = transpose(mat3x3f(view[0].xyz, view[1].xyz, view[2].xyz)); + + let T = W * J; + + let Vrk = mat3x3f( + vec3f(cov3D[0], cov3D[1], cov3D[2]), + vec3f(cov3D[1], cov3D[3], cov3D[4]), + vec3f(cov3D[2], cov3D[4], cov3D[5]) + ); + + var cov = transpose(T) * transpose(Vrk) * T; + + cov[0][0] += 0.3f; + cov[1][1] += 0.3f; + + return vec3f(cov[0][0], cov[0][1], cov[1][1]); +} + @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; + } + + let vertex = gaussians[idx]; + let a = unpack2x16float(vertex.pos_opacity[0]); + let b = unpack2x16float(vertex.pos_opacity[1]); + let pos_world = vec4f(a.x, a.y, b.x, 1.0f); + let opa = 1.0f / (1.0f + exp(-b.y)); + let pos_view = camera.view * pos_world; + let pos_clip = camera.proj * pos_view; + let pos_ndc = pos_clip.xyz / pos_clip.w; + + if (pos_ndc.x < -1.2f || pos_ndc.x > 1.2f || + pos_ndc.y < -1.2f || pos_ndc.y > 1.2f || + pos_ndc.z < 0.00f || pos_ndc.z > 1.0f) { + return; + } + + let rot_a = unpack2x16float(vertex.rot[0]); + let rot_b = unpack2x16float(vertex.rot[1]); + let rot = vec4f(rot_a.x, rot_a.y, rot_b.x, rot_b.y); + + let sca_a = unpack2x16float(vertex.scale[0]); + let sca_b = unpack2x16float(vertex.scale[1]); + let scale = exp(vec3f(sca_a.x, sca_a.y, sca_b.x)); + + let cov3D = computeCov3D(scale, scaling, rot); + let cov = computeCov2D( + pos_world, + camera.focal.x, camera.focal.y, + camera.viewport.x / (2.f * camera.focal.x), camera.viewport.y / (2.f * camera.focal.y), + cov3D, camera.view); + let det = (cov.x * cov.z - cov.y * cov.y); + + if (det < 0.000001f) { + return; + } + + let det_inv = 1.0f / det; + let conic = vec3f(cov.z * det_inv, -cov.y * det_inv, cov.x * det_inv); + + let mid = 0.5f * (cov.x + cov.z); + let lambda1 = mid + sqrt(max(0.1f, mid * mid - det)); + let lambda2 = mid - sqrt(max(0.1f, mid * mid - det)); + let radius = ceil(3.f * sqrt(max(lambda1, lambda2))); + + let keys_per_dispatch = workgroupSize * sortKeyPerThread; + let index = atomicAdd(&sort_infos.keys_size, 1u); + if (index % keys_per_dispatch == 0) { + atomicAdd(&sort_dispatch.dispatch_x, 1u); + } + + let xy = pack2x16float(pos_ndc.xy); + let wh = pack2x16float((vec2f(radius, radius) * 2.0f / camera.viewport) * scaling); + splats[index].xy = xy; + splats[index].wh = wh; + + let dir = normalize(pos_world.xyz - camera.view_inv[3].xyz); + let color = computeColorFromSH(dir, idx, shDegree); + let rg = pack2x16float(color.rg); + let ba = pack2x16float(vec2f(color.b, 1.0f)); + splats[index].rg = rg; + splats[index].ba = ba; + + let co = pack2x16float(conic.xy); + let cp = pack2x16float(vec2f(conic.z, opa)); + splats[index].co = co; + splats[index].cp = cp; - let keys_per_dispatch = workgroupSize * sortKeyPerThread; - // increment DispatchIndirect.dispatchx each time you reach limit for one dispatch of keys + sort_depths[index] = bitcast(100.0f - pos_view.z); + sort_indices[index] = index; } \ No newline at end of file