From 1f855b7bf6c73328a30c7ec9506c74652d80c572 Mon Sep 17 00:00:00 2001 From: tolelom <98kimsungmin@naver.com> Date: Thu, 26 Mar 2026 15:28:21 +0900 Subject: [PATCH] feat(renderer): add temporal accumulation compute shader for SSGI Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/voltex_renderer/src/temporal_accum.rs | 153 ++++++++++++++++++ .../voltex_renderer/src/temporal_accum.wgsl | 22 +++ 2 files changed, 175 insertions(+) create mode 100644 crates/voltex_renderer/src/temporal_accum.rs create mode 100644 crates/voltex_renderer/src/temporal_accum.wgsl diff --git a/crates/voltex_renderer/src/temporal_accum.rs b/crates/voltex_renderer/src/temporal_accum.rs new file mode 100644 index 0000000..b21f9a8 --- /dev/null +++ b/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::() 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); + } +} diff --git a/crates/voltex_renderer/src/temporal_accum.wgsl b/crates/voltex_renderer/src/temporal_accum.wgsl new file mode 100644 index 0000000..18f5ff2 --- /dev/null +++ b/crates/voltex_renderer/src/temporal_accum.wgsl @@ -0,0 +1,22 @@ +struct Params { + blend_factor: f32, + _pad: vec3, +}; + +@group(0) @binding(0) var current_tex: texture_2d; +@group(0) @binding(1) var history_tex: texture_2d; +@group(0) @binding(2) var output_tex: texture_storage_2d; +@group(0) @binding(3) var params: Params; + +@compute @workgroup_size(16, 16) +fn main(@builtin(global_invocation_id) gid: vec3) { + let dims = textureDimensions(current_tex); + if (gid.x >= dims.x || gid.y >= dims.y) { return; } + + let pos = vec2(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); +}