feat(renderer): add bilateral blur compute shader for SSGI denoising

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/voltex_renderer/src/bilateral_blur.rs

@@ -0,0 +1,156 @@
+pub struct BilateralBlur {
+    pipeline: wgpu::ComputePipeline,
+    bind_group_layout: wgpu::BindGroupLayout,
+}
+
+impl BilateralBlur {
+    pub fn new(device: &wgpu::Device) -> Self {
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("Bilateral Blur Compute"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("bilateral_blur.wgsl").into()),
+        });
+
+        let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("Bilateral Blur BGL"),
+            entries: &[
+                wgpu::BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Texture {
+                        multisampled: false,
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                        sample_type: wgpu::TextureSampleType::Float { filterable: false },
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Texture {
+                        multisampled: false,
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                        sample_type: wgpu::TextureSampleType::Depth,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 2,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Texture {
+                        multisampled: false,
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                        sample_type: wgpu::TextureSampleType::Float { filterable: false },
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 3,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::StorageTexture {
+                        access: wgpu::StorageTextureAccess::WriteOnly,
+                        format: wgpu::TextureFormat::Rgba16Float,
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                    },
+                    count: None,
+                },
+            ],
+        });
+
+        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("Bilateral Blur PL"),
+            bind_group_layouts: &[&bind_group_layout],
+            immediate_size: 0,
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Bilateral Blur Pipeline"),
+            layout: Some(&pipeline_layout),
+            module: &shader,
+            entry_point: Some("main"),
+            compilation_options: wgpu::PipelineCompilationOptions::default(),
+            cache: None,
+        });
+
+        BilateralBlur { pipeline, bind_group_layout }
+    }
+
+    pub fn dispatch(
+        &self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        input_view: &wgpu::TextureView,
+        depth_view: &wgpu::TextureView,
+        normal_view: &wgpu::TextureView,
+        output_view: &wgpu::TextureView,
+        width: u32,
+        height: u32,
+    ) {
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("Bilateral Blur BG"),
+            layout: &self.bind_group_layout,
+            entries: &[
+                wgpu::BindGroupEntry { binding: 0, resource: wgpu::BindingResource::TextureView(input_view) },
+                wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(depth_view) },
+                wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::TextureView(normal_view) },
+                wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(output_view) },
+            ],
+        });
+
+        let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("Bilateral Blur Pass"),
+            timestamp_writes: None,
+        });
+        cpass.set_pipeline(&self.pipeline);
+        cpass.set_bind_group(0, &bind_group, &[]);
+        cpass.dispatch_workgroups((width + 15) / 16, (height + 15) / 16, 1);
+    }
+}
+
+/// Pure CPU bilateral weight calculation (for testing).
+pub fn bilateral_weight(
+    spatial_dist: f32,
+    depth_diff: f32,
+    normal_dot: f32,
+    sigma_spatial: f32,
+    sigma_depth: f32,
+    sigma_normal: f32,
+) -> f32 {
+    let w_spatial = (-spatial_dist * spatial_dist / (2.0 * sigma_spatial * sigma_spatial)).exp();
+    let w_depth = (-depth_diff.abs() / sigma_depth).exp();
+    let w_normal = normal_dot.max(0.0).powf(sigma_normal);
+    w_spatial * w_depth * w_normal
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_bilateral_weight_center() {
+        // At center: dist=0, depth_diff=0, normal_dot=1
+        let w = bilateral_weight(0.0, 0.0, 1.0, 2.0, 0.1, 16.0);
+        assert!((w - 1.0).abs() < 0.01); // all factors = 1
+    }
+
+    #[test]
+    fn test_bilateral_weight_depth_edge() {
+        // Large depth difference -> low weight
+        let w = bilateral_weight(0.0, 10.0, 1.0, 2.0, 0.1, 16.0);
+        assert!(w < 0.01, "large depth diff should give low weight: {}", w);
+    }
+
+    #[test]
+    fn test_bilateral_weight_normal_edge() {
+        // Perpendicular normals -> low weight
+        let w = bilateral_weight(0.0, 0.0, 0.0, 2.0, 0.1, 16.0);
+        assert!(w < 0.01, "perpendicular normals should give low weight: {}", w);
+    }
+
+    #[test]
+    fn test_bilateral_weight_distance() {
+        // Far spatial distance -> lower weight than near
+        let w_near = bilateral_weight(1.0, 0.0, 1.0, 2.0, 0.1, 16.0);
+        let w_far = bilateral_weight(4.0, 0.0, 1.0, 2.0, 0.1, 16.0);
+        assert!(w_near > w_far);
+    }
+}

crates/voltex_renderer/src/bilateral_blur.wgsl

@@ -0,0 +1,58 @@
+@group(0) @binding(0) var input_tex: texture_2d<f32>;
+@group(0) @binding(1) var depth_tex: texture_depth_2d;
+@group(0) @binding(2) var normal_tex: texture_2d<f32>;
+@group(0) @binding(3) var output_tex: texture_storage_2d<rgba16float, write>;
+
+const KERNEL_RADIUS: i32 = 2;
+const SIGMA_SPATIAL: f32 = 2.0;
+const SIGMA_DEPTH: f32 = 0.1;
+const SIGMA_NORMAL: f32 = 16.0;
+
+fn gaussian(x: f32, sigma: f32) -> f32 {
+    return exp(-(x * x) / (2.0 * sigma * sigma));
+}
+
+@compute @workgroup_size(16, 16)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = textureDimensions(input_tex);
+    if (gid.x >= dims.x || gid.y >= dims.y) { return; }
+
+    let center = vec2<i32>(gid.xy);
+    let center_depth = textureLoad(depth_tex, center, 0);
+    let center_normal = textureLoad(normal_tex, center, 0).xyz;
+
+    var sum = vec4<f32>(0.0);
+    var weight_sum = 0.0;
+
+    for (var dy = -KERNEL_RADIUS; dy <= KERNEL_RADIUS; dy++) {
+        for (var dx = -KERNEL_RADIUS; dx <= KERNEL_RADIUS; dx++) {
+            let offset = vec2<i32>(dx, dy);
+            let sample_pos = center + offset;
+
+            if (sample_pos.x < 0 || sample_pos.y < 0 ||
+                sample_pos.x >= i32(dims.x) || sample_pos.y >= i32(dims.y)) {
+                continue;
+            }
+
+            let spatial_dist = sqrt(f32(dx * dx + dy * dy));
+            let w_spatial = gaussian(spatial_dist, SIGMA_SPATIAL);
+
+            let sample_depth = textureLoad(depth_tex, sample_pos, 0);
+            let depth_diff = abs(center_depth - sample_depth);
+            let w_depth = exp(-depth_diff / SIGMA_DEPTH);
+
+            let sample_normal = textureLoad(normal_tex, sample_pos, 0).xyz;
+            let n_dot = max(dot(center_normal, sample_normal), 0.0);
+            let w_normal = pow(n_dot, SIGMA_NORMAL);
+
+            let weight = w_spatial * w_depth * w_normal;
+            let sample_color = textureLoad(input_tex, sample_pos, 0);
+
+            sum += sample_color * weight;
+            weight_sum += weight;
+        }
+    }
+
+    let result = select(vec4<f32>(0.0), sum / weight_sum, weight_sum > 0.0);
+    textureStore(output_tex, vec2<i32>(gid.xy), result);
+}

crates/voltex_renderer/src/lib.rs

@@ -35,6 +35,8 @@ pub mod tonemap;
 pub mod forward_pass;
 pub mod auto_exposure;
 pub mod instancing;
+pub mod bilateral_blur;
+pub mod temporal_accum;
 
 pub use gpu::{GpuContext, DEPTH_FORMAT};
 pub use light::{CameraUniform, LightUniform, LightData, LightsUniform, MAX_LIGHTS, LIGHT_DIRECTIONAL, LIGHT_POINT, LIGHT_SPOT};
@@ -73,6 +75,8 @@ pub use tonemap::{TonemapUniform, aces_tonemap};
 pub use forward_pass::{ForwardPass, sort_transparent_back_to_front};
 pub use auto_exposure::AutoExposure;
 pub use instancing::{InstanceData, InstanceBuffer, create_instanced_pipeline};
+pub use bilateral_blur::BilateralBlur;
+pub use temporal_accum::TemporalAccumulation;
 pub use png::parse_png;
 pub use jpg::parse_jpg;
 pub use gltf::{parse_gltf, GltfData, GltfMesh, GltfMaterial};