feat(renderer): add temporal accumulation compute shader for SSGI

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

crates/voltex_renderer/src/temporal_accum.rs

@@ -0,0 +1,153 @@
+use bytemuck::{Pod, Zeroable};
+
+#[repr(C)]
+#[derive(Copy, Clone, Pod, Zeroable)]
+struct TemporalParams {
+    blend_factor: f32,
+    _pad: [f32; 3],
+}
+
+pub struct TemporalAccumulation {
+    pipeline: wgpu::ComputePipeline,
+    bind_group_layout: wgpu::BindGroupLayout,
+    params_buffer: wgpu::Buffer,
+    pub blend_factor: f32,
+}
+
+impl TemporalAccumulation {
+    pub fn new(device: &wgpu::Device) -> Self {
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("Temporal Accumulation Compute"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("temporal_accum.wgsl").into()),
+        });
+
+        let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("Temporal Accum 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::Float { filterable: false } },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 2, visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::StorageTexture {
+                        access: wgpu::StorageTextureAccess::WriteOnly,
+                        format: wgpu::TextureFormat::Rgba16Float,
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 3, visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
+                    count: None,
+                },
+            ],
+        });
+
+        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("Temporal Accum PL"),
+            bind_group_layouts: &[&bind_group_layout],
+            immediate_size: 0,
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Temporal Accum Pipeline"),
+            layout: Some(&pipeline_layout),
+            module: &shader,
+            entry_point: Some("main"),
+            compilation_options: wgpu::PipelineCompilationOptions::default(),
+            cache: None,
+        });
+
+        let params_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Temporal Params"),
+            size: std::mem::size_of::<TemporalParams>() as u64,
+            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        TemporalAccumulation {
+            pipeline,
+            bind_group_layout,
+            params_buffer,
+            blend_factor: 0.1,
+        }
+    }
+
+    pub fn dispatch(
+        &self,
+        device: &wgpu::Device,
+        queue: &wgpu::Queue,
+        encoder: &mut wgpu::CommandEncoder,
+        current_view: &wgpu::TextureView,
+        history_view: &wgpu::TextureView,
+        output_view: &wgpu::TextureView,
+        width: u32,
+        height: u32,
+    ) {
+        let params = TemporalParams { blend_factor: self.blend_factor, _pad: [0.0; 3] };
+        queue.write_buffer(&self.params_buffer, 0, bytemuck::cast_slice(&[params]));
+
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("Temporal Accum BG"),
+            layout: &self.bind_group_layout,
+            entries: &[
+                wgpu::BindGroupEntry { binding: 0, resource: wgpu::BindingResource::TextureView(current_view) },
+                wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(history_view) },
+                wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::TextureView(output_view) },
+                wgpu::BindGroupEntry { binding: 3, resource: self.params_buffer.as_entire_binding() },
+            ],
+        });
+
+        let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("Temporal Accum 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 temporal blend (for testing).
+pub fn temporal_blend(history: f32, current: f32, factor: f32) -> f32 {
+    history * (1.0 - factor) + current * factor
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_blend_factor_zero() {
+        let result = temporal_blend(0.5, 1.0, 0.0);
+        assert!((result - 0.5).abs() < 1e-6); // all history
+    }
+
+    #[test]
+    fn test_blend_factor_one() {
+        let result = temporal_blend(0.5, 1.0, 1.0);
+        assert!((result - 1.0).abs() < 1e-6); // all current
+    }
+
+    #[test]
+    fn test_blend_factor_half() {
+        let result = temporal_blend(0.0, 1.0, 0.5);
+        assert!((result - 0.5).abs() < 1e-6);
+    }
+
+    #[test]
+    fn test_blend_default() {
+        // Default factor = 0.1: 90% history + 10% current
+        let result = temporal_blend(0.8, 0.2, 0.1);
+        let expected = 0.8 * 0.9 + 0.2 * 0.1;
+        assert!((result - expected).abs() < 1e-6);
+    }
+}

crates/voltex_renderer/src/temporal_accum.wgsl

@@ -0,0 +1,22 @@
+struct Params {
+    blend_factor: f32,
+    _pad: vec3<f32>,
+};
+
+@group(0) @binding(0) var current_tex: texture_2d<f32>;
+@group(0) @binding(1) var history_tex: texture_2d<f32>;
+@group(0) @binding(2) var output_tex: texture_storage_2d<rgba16float, write>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+@compute @workgroup_size(16, 16)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = textureDimensions(current_tex);
+    if (gid.x >= dims.x || gid.y >= dims.y) { return; }
+
+    let pos = vec2<i32>(gid.xy);
+    let current = textureLoad(current_tex, pos, 0);
+    let history = textureLoad(history_tex, pos, 0);
+
+    let result = mix(history, current, params.blend_factor);
+    textureStore(output_tex, pos, result);
+}