diff --git a/.github/workflows/static.yml b/.github/workflows/static.yml index 9432148..46021d2 100644 --- a/.github/workflows/static.yml +++ b/.github/workflows/static.yml @@ -33,7 +33,7 @@ jobs: - name: Set up Node uses: actions/setup-node@v4 with: - node-version: 20 + node-version: 22 cache: 'npm' - name: Install dependencies run: npm ci diff --git a/.gitignore b/.gitignore index 054e565..287928e 100644 --- a/.gitignore +++ b/.gitignore @@ -24,4 +24,4 @@ dist-ssr *.sw? /.vite -*/scenes \ No newline at end of file +scenes/ \ No newline at end of file diff --git a/README.md b/README.md index edffdaf..312773a 100644 --- a/README.md +++ b/README.md @@ -2,30 +2,65 @@ **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) +* Aaron Tian + * [LinkedIn](https://www.linkedin.com/in/aaron-c-tian/), [personal website](https://aarontian-stack.github.io/) +* Tested on: Windows 22H2 (26100.6584), Intel Core Ultra 7 265k @ 3.90GHz, 32GB RAM, RTX 5070 12GB (release driver 581.15) -### Live Demo +## Live Demo -[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +[Link](https://aarontian-stack.github.io/Project5-WebGPU-Gaussian-Splat-Viewer/) -### Demo Video/GIF +[![live demo](images/demo.gif)](https://aarontian-stack.github.io/Project5-WebGPU-Gaussian-Splat-Viewer/) -[![](img/video.mp4)](TODO) +## Summary -### (TODO: Your README) +A WebGPU app for rendering Gaussian splats as described in [3D Gaussian Splatting for Real-Time Radiance Field Rendering](https://repo-sam.inria.fr/fungraph/3d-gaussian-splatting/). This allows for real time rendering of high quality radiance fields. -*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. +## Implementation -This assignment has a considerable amount of performance analysis compared -to implementation work. Complete the implementation early to leave time! +The app renders pretrained Gaussian splat models in the `.ply` format. A preprocessing compute shader is used to perform tasks such as transforming the Gaussians' positions, calculating the covariance matrix, and calculating size and color (this is the "splatting"). The Gaussians are then sorted back-to-front using radix sort. Finally the Gaussians (now splats) are rendered as quads using a single indirect draw call that had its instance count set in the preprocessing step. -### Credits +### Preprocessing Compute Shader + +Gaussian positions are transformed into NDC space using the camera's view-projection matrix. The covariance matrix is calculated from the Gaussian's rotation and scale. Each Gaussian stores spherical harmonic coefficients for calculating color based off the view direction. This information is appended to a buffer of 2D splats to be rendered along with a separate buffer for sorting those splats. This also keeps track of the number of splats that survive view-frustum culling, allowing the number of splats to be sorted to be copied to the indirect draw buffer later (the workgroup size for the sorting pass is also calculated here). + +### Rendering + +Each splat is rendered as a quad. The size simply comes from the result in the preprocessing step. For each fragment of the quad we determine if it is inside the splat (ellipse) using the centered matrix equation. If the test passes, we output the color of the splat. The opacity of the color decays exponentially based off the distance of the center of the splat. + +## Performance Analysis + +With either the Bicycle or Bonsai splat scenes, the average frame rate is the same and never drops below my max refresh rate of 100 FPS, regardless of settings, so I discuss theoretical performance implications of different settings below. + +### Point Cloud vs Gaussian Splatting + +The point cloud renders each Gaussian using point primitives, so no triangles are involved compared to splatting which uses quads (2 triangles per Gaussian). There is also no shading work in the fragment shader for point clouds (outputs constant color). Thus the Gaussian splatting incurs the additional costs of: +* Triangle rasterization +* Shading/Blending + +### Workgroup Size + +A larger workgroup size in the preprocessing step may offer benefits such as better memory coalescing, especially considering our work is a 1D dispatch. We do not use shared memory in the shader, so that is not relevant here. If the preprocessing step uses too many registers on a CU it might be better to reduce the workgroup size to mitigate spilling to shared memory. However to my knowledge it is not possible to determine this kind of information in WebGPU. + + +### View-frustum Culling + +By culling splats that are outside the view-frustum, we can reduce the number of splats that need to be sorted and rendered. This should reduce compute pressure (from sorting step) and graphics pipeline pressure (less rasterization, shading, blending). The check for culling in the preprocessing shader is very simple, but it causes an early return. Depending on the order of the Gaussians in memory, this could cause some divergence in the preprocessing step. + +### Number of Gaussians + +A larger amount of Gaussians should increase the workload in all parts of the application. Loading from disk will take longer, preprocessing will take longer, sorting will take longer, and rendering will take longer due to more primitives. + +## Bloopers + +My Gaussians got flattened into a line... + +![wtf](images/wtf.png) + +## 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) +- Special Thanks to: Shrek Shao (Google WebGPU team) & [Differential Gaussian Renderer](https://github.com/graphdeco-inria/diff-gaussian-rasterization) diff --git a/images/bike.png b/images/bike.png new file mode 100644 index 0000000..f7c80b5 Binary files /dev/null and b/images/bike.png differ diff --git a/images/bonsai.png b/images/bonsai.png new file mode 100644 index 0000000..d10b57d Binary files /dev/null and b/images/bonsai.png differ diff --git a/images/demo.gif b/images/demo.gif new file mode 100644 index 0000000..509fe94 Binary files /dev/null and b/images/demo.gif differ diff --git a/images/wtf.png b/images/wtf.png new file mode 100644 index 0000000..a315d02 Binary files /dev/null and b/images/wtf.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..f82b459 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 { - + update_scaling: (scaling: number) => void; } // Utility to create GPU buffers @@ -35,13 +35,120 @@ export default function get_renderer( // =============================================== const nulling_data = new Uint32Array([0]); + // To clear number of points to 0 before preprocess + const null_buffer = createBuffer( + device, + 'null buffer', + 4, + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + nulling_data + ); + device.queue.writeBuffer(null_buffer, 0, nulling_data); + + // vertexCount: u32, instanceCount: u32, firstVertex: u32, firstInstance: u32 + const indirect_draw_buffer = createBuffer( + device, + 'indirect draw', + 4 * 4, + GPUBufferUsage.INDIRECT | GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, + new Uint32Array([6, pc.num_points, 0, 0]) + ); + + const splat_buffer = createBuffer( + device, + 'splat buffer', + pc.num_points * 5 * 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + + const render_settings_buffer = createBuffer( + device, + 'render settings', + 8, // 2 f32 + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + new Float32Array([1.0, pc.sh_deg]) // default scaling 1.0 + ); + + // 'auto' seemingly removes unused bindings!!! So make them manually + + const camera_bind_group_layout = device.createBindGroupLayout({ + label: 'camera bind group layout', + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE | GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT, + buffer: { type: 'uniform' } + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE | GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT, + buffer: { type: 'uniform' } + } + ] + }); + + const gaussian_bind_group_layout = device.createBindGroupLayout({ + label: 'gaussian bind group layout', + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'read-only-storage' } + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'storage' } + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'read-only-storage' } + } + ] + }); + + const sort_bind_group_layout = device.createBindGroupLayout({ + label: 'sort bind group layout', + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'storage' } + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'storage' } + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'storage' } + }, + { + binding: 3, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'storage' } + } + ] + }); + + const preprocess_pipeline_layout = device.createPipelineLayout({ + label: 'preprocess pipeline layout', + bindGroupLayouts: [ + camera_bind_group_layout, + gaussian_bind_group_layout, + sort_bind_group_layout + ] + }); // =============================================== // Create Compute Pipeline and Bind Groups // =============================================== const preprocess_pipeline = device.createComputePipeline({ label: 'preprocess', - layout: 'auto', + layout: preprocess_pipeline_layout, compute: { module: device.createShaderModule({ code: preprocessWGSL }), entryPoint: 'preprocess', @@ -54,7 +161,7 @@ export default function get_renderer( const sort_bind_group = device.createBindGroup({ label: 'sort', - layout: preprocess_pipeline.getBindGroupLayout(2), + layout: sort_bind_group_layout, entries: [ { binding: 0, resource: { buffer: sorter.sort_info_buffer } }, { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_depths_buffer } }, @@ -68,19 +175,147 @@ export default function get_renderer( // Create Render Pipeline and Bind Groups // =============================================== + const render_splat_bind_group_layout = device.createBindGroupLayout({ + label: 'render splat bind group layout', + entries: [ + { + binding: 0, + visibility: GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT, + buffer: { type: 'read-only-storage' } + }, + { + binding: 1, + visibility: GPUShaderStage.VERTEX, + buffer: { type: 'read-only-storage' } + } + ] + }); + + const render_pipeline_layout = device.createPipelineLayout({ + label: 'render pipeline layout', + bindGroupLayouts: [ + camera_bind_group_layout, + render_splat_bind_group_layout + ] + }); + + const render_shader = device.createShaderModule({code: renderWGSL}); + const render_pipeline = device.createRenderPipeline({ + label: 'gaussian render', + layout: render_pipeline_layout, + vertex: { + module: render_shader, + entryPoint: 'vs_main', + }, + fragment: { + module: render_shader, + entryPoint: 'fs_main', + targets: [{ + format: presentation_format, + blend: { + color: { + srcFactor: 'src-alpha', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + alpha: { + srcFactor: 'one', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + }, + }], + }, + primitive: { + topology: 'triangle-list', + cullMode: 'none', + frontFace: 'ccw', + } + }); + + const camera_settings_bind_group = device.createBindGroup({ + label: 'gaussian camera', + layout: camera_bind_group_layout, + entries: [ + {binding: 0, resource: { buffer: camera_buffer }}, + {binding: 1, resource: { buffer: render_settings_buffer }}, + ], + }); + + const gaussian_splat_bind_group = device.createBindGroup({ + label: 'gaussian splats', + layout: gaussian_bind_group_layout, + entries: [ + { binding: 0, resource: { buffer: pc.gaussian_3d_buffer } }, + { binding: 1, resource: { buffer: splat_buffer } }, + { binding: 2, resource: { buffer: pc.sh_buffer } }, + ], + }); + + const splat_bind_group = device.createBindGroup({ + label: 'splat bind group', + layout: render_splat_bind_group_layout, + entries: [ + { binding: 0, resource: { buffer: splat_buffer } }, + { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer } }, + ], + }); // =============================================== // Command Encoder Functions // =============================================== - + let preprocess_func = (encoder: GPUCommandEncoder) => { + let preprocess_pass = encoder.beginComputePass({ + label: 'preprocess', + }); + preprocess_pass.setPipeline(preprocess_pipeline); + preprocess_pass.setBindGroup(0, camera_settings_bind_group); + preprocess_pass.setBindGroup(1, gaussian_splat_bind_group); + preprocess_pass.setBindGroup(2, sort_bind_group); + const workgroups_needed = Math.ceil(pc.num_points / C.histogram_wg_size); + preprocess_pass.dispatchWorkgroups(workgroups_needed); + preprocess_pass.end(); + }; + + let render_func = (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + const pass = encoder.beginRenderPass({ + label: 'gaussian render', + colorAttachments: [ + { + view: texture_view, + loadOp: 'clear', + storeOp: 'store', + } + ], + }); + pass.setPipeline(render_pipeline); + pass.setBindGroup(0, camera_settings_bind_group); + pass.setBindGroup(1, splat_bind_group); + + pass.drawIndirect(indirect_draw_buffer, 0); + pass.end(); + }; // =============================================== // Return Render Object // =============================================== return { frame: (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + + encoder.copyBufferToBuffer(null_buffer, 0, sorter.sort_info_buffer, 0, 4); // Clear keys_size to 0 + encoder.copyBufferToBuffer(null_buffer, 0, sorter.sort_dispatch_indirect_buffer, 0, 4); // Clear dispatch_x to 0 + + preprocess_func(encoder); + sorter.sort(encoder); + + encoder.copyBufferToBuffer(sorter.sort_info_buffer, 0, indirect_draw_buffer, 4, 4); // Copy key size to indirect draw buffer + + render_func(encoder, texture_view); }, camera_buffer, + update_scaling: (scaling: number) => { + device.queue.writeBuffer(render_settings_buffer, 0, new Float32Array([scaling, pc.sh_deg])); + }, }; } diff --git a/src/renderers/renderer.ts b/src/renderers/renderer.ts index ffdf9ba..534cf3b 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) { + gaussian_renderer.update_scaling(e.value); + } }); } diff --git a/src/shaders/gaussian.wgsl b/src/shaders/gaussian.wgsl index 759226d..5f8fa93 100644 --- a/src/shaders/gaussian.wgsl +++ b/src/shaders/gaussian.wgsl @@ -1,22 +1,85 @@ +struct CameraUniforms { + view: mat4x4, + view_inv: mat4x4, + proj: mat4x4, + proj_inv: mat4x4, + viewport: vec2, + focal: vec2 +}; + +struct RenderSettings { + gaussian_scaling: f32, + sh_deg: f32, +}; + struct VertexOutput { @builtin(position) position: vec4, - //TODO: information passed from vertex shader to fragment shader + @location(0) color: vec4, + @location(1) pixel_center: vec2, + @location(2) conic: vec3, }; struct Splat { - //TODO: information defined in preprocess compute shader + position: u32, // 2 f16 + color: array, // 4 f16 + conic: array, // 4 f16. cov + radius }; +@group(0) @binding(0) var camera: CameraUniforms; +@group(0) @binding(1) var render_settings: RenderSettings; + +@group(1) @binding(0) var splats: array; +@group(1) @binding(1) var sort_indices: array; + @vertex 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 = sort_indices[instance_index]; + let splat = splats[splat_idx]; + + let center = unpack2x16float(splat.position); + let conic_0 = unpack2x16float(splat.conic[0]); + let conic_1 = unpack2x16float(splat.conic[1]); + let quad_size = vec2(conic_1.y, conic_1.y) / camera.viewport; + + let offsets = array, 6>( + vec2(-quad_size.x, -quad_size.y), + vec2(quad_size.x, -quad_size.y), + vec2(-quad_size.x, quad_size.y), + vec2(-quad_size.x, quad_size.y), + vec2(quad_size.x, -quad_size.y), + vec2(quad_size.x, quad_size.y) + ); + out.position = vec4(center + offsets[vertex_index], 0.0, 1.0); + out.pixel_center = (center * vec2f(0.5, -0.5) + 0.5) * camera.viewport; + out.conic = vec3(conic_0.x, conic_0.y, conic_1.x); + + let color = array( + f32(unpack2x16float(splat.color[0]).x), + f32(unpack2x16float(splat.color[0]).y), + f32(unpack2x16float(splat.color[1]).x), + f32(unpack2x16float(splat.color[1]).y) + ); + out.color = vec4(color[0], color[1], color[2], color[3]); + return out; } @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4 { - return vec4(1.); + let d = in.position.xy - in.pixel_center; + + let power = -0.5 * (in.conic.x * d.x * d.x + in.conic.z * d.y * d.y) - in.conic.y * d.x * d.y; + + if (power > 0.0) { + return vec4(0.0); + } + + let alpha = min(0.99, in.color.a * exp(power)); + + return vec4(in.color.rgb, alpha); } \ 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..b50c234 100644 --- a/src/shaders/preprocess.wgsl +++ b/src/shaders/preprocess.wgsl @@ -56,23 +56,42 @@ struct Gaussian { }; struct Splat { - //TODO: store information for 2D splat rendering + position: u32, // 2 f16 + color: array, // 4 f16 + conic: array, // 4 f16. cov + radius }; -//TODO: bind your data here -@group(2) @binding(0) -var sort_infos: SortInfos; -@group(2) @binding(1) -var sort_depths : array; -@group(2) @binding(2) -var sort_indices : array; -@group(2) @binding(3) -var sort_dispatch: DispatchIndirect; +@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 sh_coeffs: array; + +@group(2) @binding(0) var sort_infos: SortInfos; +@group(2) @binding(1) var sort_depths: array; +@group(2) @binding(2) var sort_indices: array; +@group(2) @binding(3) var sort_dispatch: DispatchIndirect; + +fn fast_modulo(a: u32, b: u32) -> u32 { + return a - (a / b) * b; +} /// 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); + // 16 max coefficients, each is f16 so 8*3 for 3 channels + // Offset by coefficient, divide by 2 for f16. Then get color in that coefficient + let channel = fast_modulo(c_idx, 2); + let base_idx = splat_idx * 24 + (c_idx / 2) * 3 + channel; + + let color01 = unpack2x16float(sh_coeffs[base_idx + 0]); + let color23 = unpack2x16float(sh_coeffs[base_idx + 1]); + + if (fast_modulo(c_idx, 2) == 0u) { + return vec3f(color01.x, color01.y, color23.x); + } + + return vec3f(color01.y, color23.x, color23.y); } // spherical harmonics evaluation with Condon–Shortley phase @@ -111,8 +130,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; + } + + // Transform Gaussian position to NDC for culling + let gaussian = gaussians[idx]; + + // Unpack position and opacity + let pos_XY = unpack2x16float(gaussian.pos_opacity[0]); // x, y + let pos_Y_opacity = unpack2x16float(gaussian.pos_opacity[1]); // z, opacity + let pos = vec3(pos_XY.x, pos_XY.y, pos_Y_opacity.x); + let opacity = 1.0 / (1.0 + exp(-pos_Y_opacity.y)); // Sigmoid + + let view_pos = camera.view * vec4(pos, 1.0); + var pos_ndc = camera.proj * view_pos; + pos_ndc /= pos_ndc.w; + if (abs(pos_ndc.x) > 1.2 || abs(pos_ndc.y) > 1.2 || view_pos.z < 0.0) { + return; + } + + // Unpack rotation and scale + let rot_WX = unpack2x16float(gaussian.rot[0]); + let rot_YZ = unpack2x16float(gaussian.rot[1]); + let rot = vec4(rot_WX.y, rot_YZ.x, rot_YZ.y, rot_WX.x); // XYZW quaternion + + let scale_0 = exp(unpack2x16float(gaussian.scale[0])); // scale_0, scale_1 + let scale_1 = exp(unpack2x16float(gaussian.scale[1])); // scale_2, padding + let scale = vec3(scale_0.x, scale_0.y, scale_1.x); + + let R = mat3x3( + 1.0 - 2.0 * (rot.y * rot.y + rot.z * rot.z), 2.0 * (rot.x * rot.y - rot.w * rot.z), 2.0 * (rot.x * rot.z + rot.w * rot.y), + 2.0 * (rot.x * rot.y + rot.w * rot.z), 1.0 - 2.0 * (rot.x * rot.x + rot.z * rot.z), 2.0 * (rot.y * rot.z - rot.w * rot.x), + 2.0 * (rot.x * rot.z - rot.w * rot.y), 2.0 * (rot.y * rot.z + rot.w * rot.x), 1.0 - 2.0 * (rot.x * rot.x + rot.y * rot.y) + ); + let S = mat3x3( + render_settings.gaussian_scaling * scale.x, 0.0, 0.0, + 0.0, render_settings.gaussian_scaling * scale.y, 0.0, + 0.0, 0.0, render_settings.gaussian_scaling * scale.z + ); + + let cov3D = transpose(S * R) * S * R; + + let t = view_pos.xyz; + let J = mat3x3( + camera.focal.x / t.z, 0.0, -(camera.focal.x * t.x) / (t.z * t.z), + 0.0, camera.focal.y / t.z, -(camera.focal.y * t.y) / (t.z * t.z), + 0.0, 0.0, 0.0 + ); + + let W = transpose(mat3x3(camera.view[0].xyz, camera.view[1].xyz, camera.view[2].xyz)); + + let T = W * J; + + let Vrk = mat3x3( + 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) * transpose(Vrk) * T; + cov2D[0][0] += 0.3; + cov2D[1][1] += 0.3; + + let cov2D_flat = vec3( + cov2D[0][0], + cov2D[0][1], + cov2D[1][1] + ); + + let det = cov2D_flat.x * cov2D_flat.z - cov2D_flat.y * cov2D_flat.y; + if (det == 0.0) { + return; + } + let det_inv = 1.0 / det; + let conic = vec3(cov2D_flat.z * det_inv, -cov2D_flat.y * det_inv, cov2D_flat.x * det_inv); // Upper triangle of matrix + + let mid = 0.5 * (cov2D_flat.x + cov2D_flat.z); + let lambda1 = mid + sqrt(max(0.1, mid * mid - det)); + let lambda2 = mid - sqrt(max(0.1, mid * mid - det)); + let radius = ceil(3.0 * sqrt(max(lambda1, lambda2))); + + let cam_pos = -camera.view[3].xyz; + let view_dir = normalize(pos - cam_pos); + let color = computeColorFromSH(view_dir, idx, u32(render_settings.sh_deg)); + + // Append the splat + let splat_idx = atomicAdd(&sort_infos.keys_size, 1u); + splats[splat_idx].position = pack2x16float(pos_ndc.xy); + splats[splat_idx].color[0] = pack2x16float(color.xy); + splats[splat_idx].color[1] = pack2x16float(vec2(color.z, opacity)); + splats[splat_idx].conic[0] = pack2x16float(conic.xy); + splats[splat_idx].conic[1] = pack2x16float(vec2(conic.z, radius)); + + sort_indices[splat_idx] = splat_idx; + + var depth_uint = bitcast(-view_pos.z); + // Flip bits for sorting in radix + // If sign bit is set (negative), flip all bits + // If sign bit is not set (positive), flip only sign bit + let mask = select(0x80000000u, 0xFFFFFFFFu, (depth_uint & 0x80000000u) != 0u); + sort_depths[splat_idx] = depth_uint ^ mask; 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 == 0u) { + atomicAdd(&sort_dispatch.dispatch_x, 1u); + } } \ No newline at end of file