From 447473598a075a4e43f2dadce3382bfcbba03ee5 Mon Sep 17 00:00:00 2001 From: tolelom <98kimsungmin@naver.com> Date: Thu, 26 Mar 2026 15:57:05 +0900 Subject: [PATCH] feat(renderer): add motion blur compute shader Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/voltex_renderer/src/dof.rs | 140 ++++++++++++++++ crates/voltex_renderer/src/dof.wgsl | 55 ++++++ crates/voltex_renderer/src/lib.rs | 4 + crates/voltex_renderer/src/motion_blur.rs | 177 ++++++++++++++++++++ crates/voltex_renderer/src/motion_blur.wgsl | 49 ++++++ 5 files changed, 425 insertions(+) create mode 100644 crates/voltex_renderer/src/dof.rs create mode 100644 crates/voltex_renderer/src/dof.wgsl create mode 100644 crates/voltex_renderer/src/motion_blur.rs create mode 100644 crates/voltex_renderer/src/motion_blur.wgsl diff --git a/crates/voltex_renderer/src/dof.rs b/crates/voltex_renderer/src/dof.rs new file mode 100644 index 0000000..375ce4b --- /dev/null +++ b/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::() 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); + } +} diff --git a/crates/voltex_renderer/src/dof.wgsl b/crates/voltex_renderer/src/dof.wgsl new file mode 100644 index 0000000..a3dc72b --- /dev/null +++ b/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; +@group(0) @binding(1) var depth_tex: texture_depth_2d; +@group(0) @binding(2) var output_tex: texture_storage_2d; +@group(0) @binding(3) var 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) { + let dims = textureDimensions(color_tex); + if (gid.x >= dims.x || gid.y >= dims.y) { return; } + + let pos = vec2(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(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(dx, dy); + let clamped = clamp(sample_pos, vec2(0), vec2(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); +} diff --git a/crates/voltex_renderer/src/lib.rs b/crates/voltex_renderer/src/lib.rs index 67ea82f..9b9e6d4 100644 --- a/crates/voltex_renderer/src/lib.rs +++ b/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; diff --git a/crates/voltex_renderer/src/motion_blur.rs b/crates/voltex_renderer/src/motion_blur.rs new file mode 100644 index 0000000..072f7e2 --- /dev/null +++ b/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::() 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::(); + assert_eq!(size % 16, 0); + } +} diff --git a/crates/voltex_renderer/src/motion_blur.wgsl b/crates/voltex_renderer/src/motion_blur.wgsl new file mode 100644 index 0000000..4df42b4 --- /dev/null +++ b/crates/voltex_renderer/src/motion_blur.wgsl @@ -0,0 +1,49 @@ +struct MotionBlurParams { + inv_view_proj: mat4x4, + prev_view_proj: mat4x4, + num_samples: u32, + strength: f32, + _pad: vec2, +}; + +@group(0) @binding(0) var color_tex: texture_2d; +@group(0) @binding(1) var depth_tex: texture_depth_2d; +@group(0) @binding(2) var output_tex: texture_storage_2d; +@group(0) @binding(3) var params: MotionBlurParams; + +@compute @workgroup_size(16, 16) +fn main(@builtin(global_invocation_id) gid: vec3) { + let dims = textureDimensions(color_tex); + if (gid.x >= dims.x || gid.y >= dims.y) { return; } + + let pos = vec2(gid.xy); + let uv = (vec2(gid.xy) + 0.5) / vec2(dims); + + // Reconstruct world position from depth + let depth = textureLoad(depth_tex, pos, 0); + let ndc = vec4(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(world_pos, 1.0); + let prev_ndc = prev_clip.xyz / prev_clip.w; + let prev_uv = vec2(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(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(sample_uv * vec2(dims)); + let clamped = clamp(sample_pos, vec2(0), vec2(dims) - 1); + color += textureLoad(color_tex, clamped, 0); + } + color /= f32(n); + + textureStore(output_tex, pos, color); +}