feat(renderer): add motion blur compute shader

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

crates/voltex_renderer/src/dof.rs

@@ -0,0 +1,140 @@
+use bytemuck::{Pod, Zeroable};
+
+#[repr(C)]
+#[derive(Copy, Clone, Pod, Zeroable)]
+pub struct DofParams {
+    pub focus_distance: f32,
+    pub focus_range: f32,
+    pub max_blur: f32,
+    pub _pad: f32,
+}
+
+impl DofParams {
+    pub fn new() -> Self {
+        DofParams { focus_distance: 5.0, focus_range: 3.0, max_blur: 5.0, _pad: 0.0 }
+    }
+}
+
+pub struct DepthOfField {
+    pipeline: wgpu::ComputePipeline,
+    bind_group_layout: wgpu::BindGroupLayout,
+    params_buffer: wgpu::Buffer,
+    pub enabled: bool,
+}
+
+impl DepthOfField {
+    pub fn new(device: &wgpu::Device) -> Self {
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("DOF Compute"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("dof.wgsl").into()),
+        });
+
+        let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("DOF 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::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("DOF PL"), bind_group_layouts: &[&bind_group_layout], immediate_size: 0,
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("DOF 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("DOF Params"),
+            size: std::mem::size_of::<DofParams>() as u64,
+            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        DepthOfField { pipeline, bind_group_layout, params_buffer, enabled: true }
+    }
+
+    pub fn dispatch(
+        &self, device: &wgpu::Device, queue: &wgpu::Queue, encoder: &mut wgpu::CommandEncoder,
+        color_view: &wgpu::TextureView, depth_view: &wgpu::TextureView,
+        output_view: &wgpu::TextureView, params: &DofParams, width: u32, height: u32,
+    ) {
+        queue.write_buffer(&self.params_buffer, 0, bytemuck::cast_slice(&[*params]));
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("DOF BG"), layout: &self.bind_group_layout,
+            entries: &[
+                wgpu::BindGroupEntry { binding: 0, resource: wgpu::BindingResource::TextureView(color_view) },
+                wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(depth_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("DOF 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);
+    }
+}
+
+/// CPU circle-of-confusion calculation (for testing).
+pub fn circle_of_confusion(depth: f32, focus_distance: f32, focus_range: f32, max_blur: f32) -> f32 {
+    let diff = (depth - focus_distance).abs();
+    (diff / focus_range).clamp(0.0, 1.0) * max_blur
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_coc_in_focus() {
+        let coc = circle_of_confusion(5.0, 5.0, 3.0, 5.0);
+        assert!((coc - 0.0).abs() < 1e-6); // at focus distance
+    }
+
+    #[test]
+    fn test_coc_near() {
+        let coc = circle_of_confusion(2.0, 5.0, 3.0, 5.0);
+        assert!((coc - 5.0).abs() < 1e-6); // diff=3, range=3 -> full blur
+    }
+
+    #[test]
+    fn test_coc_far() {
+        let coc = circle_of_confusion(8.0, 5.0, 3.0, 5.0);
+        assert!((coc - 5.0).abs() < 1e-6); // diff=3, range=3 -> full blur
+    }
+
+    #[test]
+    fn test_coc_partial() {
+        let coc = circle_of_confusion(6.5, 5.0, 3.0, 5.0);
+        assert!((coc - 2.5).abs() < 1e-6); // diff=1.5, range=3 -> 0.5 * 5 = 2.5
+    }
+
+    #[test]
+    fn test_params_default() {
+        let p = DofParams::new();
+        assert!((p.focus_distance - 5.0).abs() < 1e-6);
+        assert!((p.focus_range - 3.0).abs() < 1e-6);
+    }
+}

crates/voltex_renderer/src/dof.wgsl

@@ -0,0 +1,55 @@
+struct DofParams {
+    focus_distance: f32,
+    focus_range: f32,
+    max_blur: f32,
+    _pad: f32,
+};
+
+@group(0) @binding(0) var color_tex: texture_2d<f32>;
+@group(0) @binding(1) var depth_tex: texture_depth_2d;
+@group(0) @binding(2) var output_tex: texture_storage_2d<rgba16float, write>;
+@group(0) @binding(3) var<uniform> params: DofParams;
+
+fn circle_of_confusion(depth: f32) -> f32 {
+    let diff = abs(depth - params.focus_distance);
+    let coc = clamp(diff / params.focus_range, 0.0, 1.0) * params.max_blur;
+    return coc;
+}
+
+@compute @workgroup_size(16, 16)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = textureDimensions(color_tex);
+    if (gid.x >= dims.x || gid.y >= dims.y) { return; }
+
+    let pos = vec2<i32>(gid.xy);
+    let depth = textureLoad(depth_tex, pos, 0);
+    let coc = circle_of_confusion(depth);
+
+    if (coc < 0.5) {
+        // In focus — no blur
+        textureStore(output_tex, pos, textureLoad(color_tex, pos, 0));
+        return;
+    }
+
+    // Disc blur with radius = coc
+    let radius = i32(coc);
+    var color = vec4<f32>(0.0);
+    var weight = 0.0;
+
+    for (var dy = -radius; dy <= radius; dy++) {
+        for (var dx = -radius; dx <= radius; dx++) {
+            let dist = sqrt(f32(dx * dx + dy * dy));
+            if (dist > coc) { continue; }
+
+            let sample_pos = pos + vec2<i32>(dx, dy);
+            let clamped = clamp(sample_pos, vec2<i32>(0), vec2<i32>(dims) - 1);
+            let sample_color = textureLoad(color_tex, clamped, 0);
+            let w = 1.0 - dist / (coc + 0.001);
+            color += sample_color * w;
+            weight += w;
+        }
+    }
+
+    let result = select(textureLoad(color_tex, pos, 0), color / weight, weight > 0.0);
+    textureStore(output_tex, pos, result);
+}

crates/voltex_renderer/src/lib.rs

@@ -39,7 +39,11 @@ pub mod bilateral_blur;
 pub mod temporal_accum;
 pub mod taa;
 pub mod ssr;
+pub mod motion_blur;
+pub mod dof;
 
+pub use motion_blur::MotionBlur;
+pub use dof::DepthOfField;
 pub use gpu::{GpuContext, DEPTH_FORMAT};
 pub use light::{CameraUniform, LightUniform, LightData, LightsUniform, MAX_LIGHTS, LIGHT_DIRECTIONAL, LIGHT_POINT, LIGHT_SPOT};
 pub use mesh::Mesh;

crates/voltex_renderer/src/motion_blur.rs

@@ -0,0 +1,177 @@
+use bytemuck::{Pod, Zeroable};
+
+#[repr(C)]
+#[derive(Copy, Clone, Pod, Zeroable)]
+pub struct MotionBlurParams {
+    pub inv_view_proj: [[f32; 4]; 4],
+    pub prev_view_proj: [[f32; 4]; 4],
+    pub num_samples: u32,
+    pub strength: f32,
+    pub _pad: [f32; 2],
+}
+
+impl MotionBlurParams {
+    pub fn new() -> Self {
+        MotionBlurParams {
+            inv_view_proj: [[0.0; 4]; 4],
+            prev_view_proj: [[0.0; 4]; 4],
+            num_samples: 8,
+            strength: 1.0,
+            _pad: [0.0; 2],
+        }
+    }
+}
+
+pub struct MotionBlur {
+    pipeline: wgpu::ComputePipeline,
+    bind_group_layout: wgpu::BindGroupLayout,
+    params_buffer: wgpu::Buffer,
+    pub enabled: bool,
+}
+
+impl MotionBlur {
+    pub fn new(device: &wgpu::Device) -> Self {
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("Motion Blur Compute"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("motion_blur.wgsl").into()),
+        });
+
+        let bind_group_layout =
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("Motion 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::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("Motion Blur PL"),
+            bind_group_layouts: &[&bind_group_layout],
+            immediate_size: 0,
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Motion Blur 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("Motion Blur Params"),
+            size: std::mem::size_of::<MotionBlurParams>() as u64,
+            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        MotionBlur {
+            pipeline,
+            bind_group_layout,
+            params_buffer,
+            enabled: true,
+        }
+    }
+
+    pub fn dispatch(
+        &self,
+        device: &wgpu::Device,
+        queue: &wgpu::Queue,
+        encoder: &mut wgpu::CommandEncoder,
+        color_view: &wgpu::TextureView,
+        depth_view: &wgpu::TextureView,
+        output_view: &wgpu::TextureView,
+        params: &MotionBlurParams,
+        width: u32,
+        height: u32,
+    ) {
+        queue.write_buffer(&self.params_buffer, 0, bytemuck::cast_slice(&[*params]));
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("Motion Blur BG"),
+            layout: &self.bind_group_layout,
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(color_view),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::TextureView(depth_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("Motion 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);
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_params_default() {
+        let p = MotionBlurParams::new();
+        assert_eq!(p.num_samples, 8);
+        assert!((p.strength - 1.0).abs() < 1e-6);
+    }
+
+    #[test]
+    fn test_params_size_aligned() {
+        let size = std::mem::size_of::<MotionBlurParams>();
+        assert_eq!(size % 16, 0);
+    }
+}

crates/voltex_renderer/src/motion_blur.wgsl

@@ -0,0 +1,49 @@
+struct MotionBlurParams {
+    inv_view_proj: mat4x4<f32>,
+    prev_view_proj: mat4x4<f32>,
+    num_samples: u32,
+    strength: f32,
+    _pad: vec2<f32>,
+};
+
+@group(0) @binding(0) var color_tex: texture_2d<f32>;
+@group(0) @binding(1) var depth_tex: texture_depth_2d;
+@group(0) @binding(2) var output_tex: texture_storage_2d<rgba16float, write>;
+@group(0) @binding(3) var<uniform> params: MotionBlurParams;
+
+@compute @workgroup_size(16, 16)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = textureDimensions(color_tex);
+    if (gid.x >= dims.x || gid.y >= dims.y) { return; }
+
+    let pos = vec2<i32>(gid.xy);
+    let uv = (vec2<f32>(gid.xy) + 0.5) / vec2<f32>(dims);
+
+    // Reconstruct world position from depth
+    let depth = textureLoad(depth_tex, pos, 0);
+    let ndc = vec4<f32>(uv.x * 2.0 - 1.0, 1.0 - uv.y * 2.0, depth, 1.0);
+    let world = params.inv_view_proj * ndc;
+    let world_pos = world.xyz / world.w;
+
+    // Project to previous frame
+    let prev_clip = params.prev_view_proj * vec4<f32>(world_pos, 1.0);
+    let prev_ndc = prev_clip.xyz / prev_clip.w;
+    let prev_uv = vec2<f32>(prev_ndc.x * 0.5 + 0.5, 1.0 - (prev_ndc.y * 0.5 + 0.5));
+
+    // Velocity = current_uv - prev_uv
+    let velocity = (uv - prev_uv) * params.strength;
+
+    // Sample along velocity direction
+    var color = vec4<f32>(0.0);
+    let n = params.num_samples;
+    for (var i = 0u; i < n; i++) {
+        let t = f32(i) / f32(n - 1u) - 0.5; // -0.5 to 0.5
+        let sample_uv = uv + velocity * t;
+        let sample_pos = vec2<i32>(sample_uv * vec2<f32>(dims));
+        let clamped = clamp(sample_pos, vec2<i32>(0), vec2<i32>(dims) - 1);
+        color += textureLoad(color_tex, clamped, 0);
+    }
+    color /= f32(n);
+
+    textureStore(output_tex, pos, color);
+}