From 1ea2d340e6f8f5b2c5cc95028b082c28e3391aa9 Mon Sep 17 00:00:00 2001 From: tolelom <98kimsungmin@naver.com> Date: Thu, 26 Mar 2026 15:28:20 +0900 Subject: [PATCH] feat(renderer): add bilateral blur compute shader for SSGI denoising Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/voltex_renderer/src/bilateral_blur.rs | 156 ++++++++++++++++++ .../voltex_renderer/src/bilateral_blur.wgsl | 58 +++++++ crates/voltex_renderer/src/lib.rs | 4 + 3 files changed, 218 insertions(+) create mode 100644 crates/voltex_renderer/src/bilateral_blur.rs create mode 100644 crates/voltex_renderer/src/bilateral_blur.wgsl diff --git a/crates/voltex_renderer/src/bilateral_blur.rs b/crates/voltex_renderer/src/bilateral_blur.rs new file mode 100644 index 0000000..b82bfe5 --- /dev/null +++ b/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); + } +} diff --git a/crates/voltex_renderer/src/bilateral_blur.wgsl b/crates/voltex_renderer/src/bilateral_blur.wgsl new file mode 100644 index 0000000..f7b92a3 --- /dev/null +++ b/crates/voltex_renderer/src/bilateral_blur.wgsl @@ -0,0 +1,58 @@ +@group(0) @binding(0) var input_tex: texture_2d; +@group(0) @binding(1) var depth_tex: texture_depth_2d; +@group(0) @binding(2) var normal_tex: texture_2d; +@group(0) @binding(3) var output_tex: texture_storage_2d; + +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) { + let dims = textureDimensions(input_tex); + if (gid.x >= dims.x || gid.y >= dims.y) { return; } + + let center = vec2(gid.xy); + let center_depth = textureLoad(depth_tex, center, 0); + let center_normal = textureLoad(normal_tex, center, 0).xyz; + + var sum = vec4(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(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(0.0), sum / weight_sum, weight_sum > 0.0); + textureStore(output_tex, vec2(gid.xy), result); +} diff --git a/crates/voltex_renderer/src/lib.rs b/crates/voltex_renderer/src/lib.rs index d8ce101..b3e8228 100644 --- a/crates/voltex_renderer/src/lib.rs +++ b/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};