From bbfd435154574e2a9f35b063633ba50317c80a26 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Sun, 29 Sep 2024 23:40:36 -0700 Subject: [PATCH] WIP separate spatial pass --- blade-helpers/src/hud.rs | 1 + blade-helpers/src/lib.rs | 1 + blade-render/code/ray-trace.wgsl | 162 +++++++++++++++++++++---------- blade-render/src/render/mod.rs | 134 +++++++++++++++++-------- 4 files changed, 205 insertions(+), 93 deletions(-) diff --git a/blade-helpers/src/hud.rs b/blade-helpers/src/hud.rs index 880b745f..6273083b 100644 --- a/blade-helpers/src/hud.rs +++ b/blade-helpers/src/hud.rs @@ -17,6 +17,7 @@ impl ExposeHud for blade_render::RayConfig { ui.add( egui::widgets::Slider::new(&mut self.temporal_history, 0..=50).text("Temporal history"), ); + ui.checkbox(&mut self.spatial_pass, "Spatial pass"); ui.add(egui::widgets::Slider::new(&mut self.spatial_taps, 0..=10).text("Spatial taps")); ui.add( egui::widgets::Slider::new(&mut self.spatial_tap_history, 0..=50) diff --git a/blade-helpers/src/lib.rs b/blade-helpers/src/lib.rs index db7f2ec8..bf42be0e 100644 --- a/blade-helpers/src/lib.rs +++ b/blade-helpers/src/lib.rs @@ -13,6 +13,7 @@ pub fn default_ray_config() -> blade_render::RayConfig { environment_importance_sampling: false, temporal_tap: true, temporal_history: 10, + spatial_pass: true, spatial_taps: 1, spatial_tap_history: 10, spatial_radius: 20, diff --git a/blade-render/code/ray-trace.wgsl b/blade-render/code/ray-trace.wgsl index 3e329783..7dc97b94 100644 --- a/blade-render/code/ray-trace.wgsl +++ b/blade-render/code/ray-trace.wgsl @@ -14,7 +14,7 @@ const PI: f32 = 3.1415926; const MAX_RESERVOIRS: u32 = 2u; // See "DECOUPLING SHADING AND REUSE" in // "Rearchitecting Spatiotemporal Resampling for Production" -const DECOUPLED_SHADING: bool = true; +const DECOUPLED_SHADING: bool = false; // We are considering 2x2 grid, so must be <= 4 const FACTOR_TEMPORAL_CANDIDATES: u32 = 1u; @@ -117,12 +117,12 @@ fn merge_reservoir(r: ptr, other: LiveReservoir, random return false; } } -fn unpack_reservoir(f: StoredReservoir, max_history: u32) -> LiveReservoir { +fn unpack_reservoir(f: StoredReservoir, max_history: u32, radiance: vec3) -> LiveReservoir { var r: LiveReservoir; r.selected_light_index = f.light_index; r.selected_uv = f.light_uv; r.selected_target_score = f.target_score; - r.radiance = vec3(0.0); // to be continued... + r.radiance = radiance; let history = min(f.confidence, f32(max_history)); r.weight_sum = f.contribution_weight * f.target_score * history; r.history = history; @@ -150,6 +150,7 @@ var t_flat_normal: texture_2d; var t_prev_flat_normal: texture_2d; var t_motion: texture_2d; var out_diffuse: texture_storage_2d; +var in_diffuse: texture_2d; var out_debug: texture_storage_2d; fn sample_circle(random: f32) -> vec2 { @@ -322,8 +323,9 @@ fn evaluate_sample(ls: LightSample, surface: Surface, start_pos: vec3, debu return brdf; } -fn produce_canonical(surface: Surface, position: vec3, rng: ptr, debug_len: f32) -> LiveReservoir { +fn produce_canonical(surface: Surface, position: vec3, rng: ptr, enable_debug: bool) -> LiveReservoir { var reservoir = LiveReservoir(); + let debug_len = select(0.0, surface.depth * 0.2, enable_debug); for (var i = 0u; i < parameters.num_environment_samples; i += 1u) { var ls: LightSample; if (parameters.environment_importance_sampling != 0u) { @@ -344,23 +346,17 @@ fn produce_canonical(surface: Surface, position: vec3, rng: ptr, - temporal_index: u32, count: u32, } -fn gather_neighborhood( - surface: Surface, pixel: vec2, rng: ptr, enable_debug: bool +fn gather_neighborhood_temporal( + surface: Surface, position: vec3, pixel: vec2, rng: ptr ) -> Neighborhood { if (surface.depth == 0.0) { return Neighborhood(); } - let ray_dir = get_ray_direction(camera, pixel); - let debug_len = select(0.0, surface.depth * 0.2, enable_debug); - let position = camera.position + surface.depth * ray_dir; - let center_coord = get_prev_pixel(pixel, position); let center_pixel = vec2(center_coord); // Trick to start with closer pixels: we derive the "further" @@ -369,28 +365,58 @@ fn gather_neighborhood( // First, gather the list of reservoirs to merge with var nh = Neighborhood(); - nh.canonical = produce_canonical(surface, position, rng, debug_len); - nh.temporal_index = ~0u; - let num_temporal_candidates = parameters.temporal_tap * FACTOR_TEMPORAL_CANDIDATES; - let num_candidates = num_temporal_candidates + parameters.spatial_taps * FACTOR_SPATIAL_CANDIDATES; - let max_samples = min(MAX_RESERVOIRS, 1u + parameters.spatial_taps); + let num_candidates = parameters.temporal_tap * FACTOR_TEMPORAL_CANDIDATES; + + for (var tap = 0u; tap <= num_candidates && nh.count == 0u; tap += 1u) { + let mask = vec2(tap) & vec2(1u, 2u); + let other_pixel = select(center_pixel, further_pixel, mask != vec2(0u)); + + let other_index = get_reservoir_index(other_pixel, prev_camera); + if (other_index < 0) { + continue; + } + if (prev_reservoirs[other_index].confidence == 0.0) { + continue; + } + + let other_surface = read_prev_surface(other_pixel); + let compatibility = compare_surfaces(surface, other_surface); + if (compatibility < 0.1) { + // if the surfaces are too different, there is no trust in this sample + continue; + } + + nh.reservoir_indices[0] = other_index; + nh.count = 1u; + } + + if (WRITE_DEBUG_IMAGE && debug.view_mode == DebugMode_TemporalReuse) { + var color = vec4(f32(nh.count)); + textureStore(out_debug, pixel, color); + } + + return nh; +} + +fn gather_neighborhood_spatial( + surface: Surface, pixel: vec2, rng: ptr +) -> Neighborhood { + if (surface.depth == 0.0) { + return Neighborhood(); + } + + // First, gather the list of reservoirs to merge with + var nh = Neighborhood(); + let num_candidates = parameters.spatial_taps * FACTOR_SPATIAL_CANDIDATES; + let max_samples = min(MAX_RESERVOIRS, parameters.spatial_taps); for (var tap = 0u; tap <= num_candidates && nh.count < max_samples; tap += 1u) { - var other_pixel = center_pixel; - if (tap < num_temporal_candidates) { - if (nh.temporal_index < tap) { - continue; - } - let mask = vec2(tap) & vec2(1u, 2u); - other_pixel = select(center_pixel, further_pixel, mask != vec2(0u)); - } else { - let r0 = max(center_pixel - vec2(parameters.spatial_radius), vec2(0)); - let r1 = min(center_pixel + vec2(parameters.spatial_radius + 1), vec2(prev_camera.target_size)); - other_pixel = vec2(mix(vec2(r0), vec2(r1), vec2(random_gen(rng), random_gen(rng)))); - let diff = other_pixel - center_pixel; - if (dot(diff, diff) < MIN_SPATIAL_REUSE_DISTANCE) { - continue; - } + let r0 = max(pixel - vec2(parameters.spatial_radius), vec2(0)); + let r1 = min(pixel + vec2(parameters.spatial_radius + 1), vec2(prev_camera.target_size)); + let other_pixel = vec2(mix(vec2(r0), vec2(r1), vec2(random_gen(rng), random_gen(rng)))); + let diff = other_pixel - pixel; + if (dot(diff, diff) < MIN_SPATIAL_REUSE_DISTANCE) { + continue; } let other_index = get_reservoir_index(other_pixel, prev_camera); @@ -408,14 +434,11 @@ fn gather_neighborhood( continue; } - if (tap < num_temporal_candidates) { - nh.temporal_index = nh.count; - } nh.reservoir_indices[nh.count] = other_index; nh.count += 1u; } - if (WRITE_DEBUG_IMAGE && debug.view_mode == DebugMode_SampleReuse) { + if (WRITE_DEBUG_IMAGE && debug.view_mode == DebugMode_SpatialReuse) { var color = vec4(0.0); for (var i = 0u; i < min(3u, nh.count); i += 1u) { color[i] = 1.0; @@ -436,7 +459,8 @@ fn ratio(a: f32, b: f32) -> f32 { } fn compute_restir( - surface: Surface, ray_dir: vec3, nh: Neighborhood, + surface: Surface, ray_dir: vec3, + canonical: LiveReservoir, nh: Neighborhood, max_history: u32, rng: ptr, enable_debug: bool, ) -> RestirOutput { if (surface.depth == 0.0) { @@ -445,8 +469,9 @@ fn compute_restir( return ro; } + let enable_restir_debug = (debug.draw_flags & DebugDrawFlags_RESTIR) != 0u && enable_debug; let position = camera.position + surface.depth * ray_dir; - let debug_len = select(0.0, surface.depth * 0.2, enable_debug); + let debug_len = select(0.0, surface.depth * 0.2, enable_restir_debug); var accepted_reservoir_indices = nh.reservoir_indices; // evaluate the MIS of each of the samples versus the canonical one. var reservoir = LiveReservoir(); @@ -459,7 +484,6 @@ fn compute_restir( let neighbor_index = accepted_reservoir_indices[rid]; let neighbor = prev_reservoirs[neighbor_index]; - let max_history = select(parameters.spatial_tap_history, parameters.temporal_history, rid == nh.temporal_index); var other: LiveReservoir; if (parameters.use_pairwise_mis != 0u) { let neighbor_pixel = get_pixel_from_reservoir_index(neighbor_index, prev_camera); @@ -470,14 +494,14 @@ fn compute_restir( let neighbor_position = prev_camera.position + neighbor_surface.depth * neighbor_dir; let t_canonical_at_neighbor = estimate_target_score_with_occlusion( - neighbor_surface, neighbor_position, nh.canonical.selected_light_index, nh.canonical.selected_uv, prev_acc_struct, debug_len); - let r_canonical = ratio(nh.canonical.history * nh.canonical.selected_target_score * inv_count, neighbor_history * t_canonical_at_neighbor.score); + neighbor_surface, neighbor_position, canonical.selected_light_index, canonical.selected_uv, prev_acc_struct, debug_len); + let r_canonical = ratio(canonical.history * canonical.selected_target_score * inv_count, neighbor_history * t_canonical_at_neighbor.score); mis_canonical += mis_scale * r_canonical; } let t_neighbor_at_canonical = estimate_target_score_with_occlusion( surface, position, neighbor.light_index, neighbor.light_uv, acc_struct, debug_len); - let r_neighbor = ratio(neighbor_history * neighbor.target_score, nh.canonical.history * t_neighbor_at_canonical.score * inv_count); + let r_neighbor = ratio(neighbor_history * neighbor.target_score, canonical.history * t_neighbor_at_canonical.score * inv_count); let mis_neighbor = mis_scale * r_neighbor; other.history = neighbor_history; @@ -487,8 +511,8 @@ fn compute_restir( other.weight_sum = t_neighbor_at_canonical.score * neighbor.contribution_weight * mis_neighbor; other.radiance = t_neighbor_at_canonical.color; } else { - other = unpack_reservoir(neighbor, max_history); - other.radiance = evaluate_reflected_light(surface, other.selected_light_index, other.selected_uv); + let radiance = evaluate_reflected_light(surface, neighbor.light_index, neighbor.light_uv); + other = unpack_reservoir(neighbor, max_history, radiance); } if (DECOUPLED_SHADING) { @@ -502,15 +526,15 @@ fn compute_restir( } // Finally, merge in the canonical sample - var canonical = nh.canonical; + var canonical_mod = canonical; if (parameters.use_pairwise_mis != 0) { - canonical.weight_sum *= mis_canonical / canonical.history; + canonical_mod.weight_sum *= mis_canonical / canonical.history; } if (DECOUPLED_SHADING) { - let cw = canonical.weight_sum / max(canonical.selected_target_score, 0.1); - color_and_weight += canonical.weight_sum * vec4(cw * canonical.radiance, 1.0); + let cw = canonical_mod.weight_sum / max(canonical_mod.selected_target_score, 0.1); + color_and_weight += canonical_mod.weight_sum * vec4(cw * canonical_mod.radiance, 1.0); } - merge_reservoir(&reservoir, canonical, random_gen(rng)); + merge_reservoir(&reservoir, canonical_mod, random_gen(rng)); let effective_history = select(reservoir.history, 1.0, parameters.use_pairwise_mis != 0); var ro = RestirOutput(); @@ -535,19 +559,51 @@ fn main(@builtin(global_invocation_id) global_id: vec3) { let surface = read_surface(vec2(global_id.xy)); let pixel = vec2(global_id.xy); let enable_debug = DEBUG_MODE && all(global_id.xy == debug.mouse_pos); - let enable_restir_debug = (debug.draw_flags & DebugDrawFlags_RESTIR) != 0u && enable_debug; - let neighborhood = gather_neighborhood(surface, pixel, &rng, enable_restir_debug); let ray_dir = get_ray_direction(camera, pixel); - let ro = compute_restir(surface, ray_dir, neighborhood, &rng, enable_restir_debug); + let position = camera.position + surface.depth * ray_dir; + let neighborhood = gather_neighborhood_temporal(surface, position, pixel, &rng); + let canonical = produce_canonical(surface, position, &rng, enable_debug); + let ro = compute_restir(surface, ray_dir, canonical, neighborhood, parameters.temporal_history, &rng, enable_debug); let pixel_index = get_reservoir_index(pixel, camera); reservoirs[pixel_index] = ro.reservoir; + textureStore(out_diffuse, pixel, vec4(ro.radiance, 1.0)); +} + +@compute @workgroup_size(8, 4) +fn main_spatial(@builtin(global_invocation_id) global_id: vec3) { + if (any(global_id.xy >= camera.target_size)) { + return; + } + + let global_index = global_id.y * camera.target_size.x + global_id.x; + var rng = random_init(global_index, parameters.frame_index * 2u); + + let surface = read_surface(vec2(global_id.xy)); + let pixel = vec2(global_id.xy); + let pixel_index = get_reservoir_index(pixel, camera); + let in_radiance = textureLoad(in_diffuse, pixel, 0).xyz; + + if (surface.depth == 0.0) { + reservoirs[pixel_index] = prev_reservoirs[pixel_index]; + textureStore(out_diffuse, pixel, vec4(in_radiance, 1.0)); + return; + } + + let enable_debug = DEBUG_MODE && all(global_id.xy == debug.mouse_pos); + + let neighborhood = gather_neighborhood_spatial(surface, pixel, &rng); + let canonical = unpack_reservoir(prev_reservoirs[pixel_index], ~0u, in_radiance); + let ray_dir = get_ray_direction(camera, pixel); + let ro = compute_restir(surface, ray_dir, canonical, neighborhood, parameters.spatial_tap_history, &rng, enable_debug); + reservoirs[pixel_index] = ro.reservoir; let color = ro.radiance; + if (enable_debug) { debug_buf.variance.color_sum += color; debug_buf.variance.color2_sum += color * color; debug_buf.variance.count += 1u; } - textureStore(out_diffuse, global_id.xy, vec4(color, 1.0)); + textureStore(out_diffuse, pixel, vec4(color, 1.0)); } diff --git a/blade-render/src/render/mod.rs b/blade-render/src/render/mod.rs index 4bbedf21..33bb7314 100644 --- a/blade-render/src/render/mod.rs +++ b/blade-render/src/render/mod.rs @@ -56,7 +56,8 @@ pub enum DebugMode { ShadingNormal = 7, Motion = 8, HitConsistency = 9, - SampleReuse = 10, + TemporalReuse = 10, + SpatialReuse = 11, Variance = 15, } @@ -97,6 +98,7 @@ pub struct RayConfig { pub environment_importance_sampling: bool, pub temporal_tap: bool, pub temporal_history: u32, + pub spatial_pass: bool, pub spatial_taps: u32, pub spatial_tap_history: u32, pub spatial_radius: u32, @@ -319,9 +321,9 @@ struct Blur { pub struct Renderer { shaders: Shaders, targets: RestirTargets, - post_proc_input_index: usize, fill_pipeline: blade_graphics::ComputePipeline, main_pipeline: blade_graphics::ComputePipeline, + spatial_pipeline: blade_graphics::ComputePipeline, post_proc_pipeline: blade_graphics::RenderPipeline, blur: Blur, acceleration_structure: blade_graphics::AccelerationStructure, @@ -426,6 +428,7 @@ struct MainData { reservoirs: blade_graphics::BufferPiece, prev_reservoirs: blade_graphics::BufferPiece, out_diffuse: blade_graphics::TextureView, + in_diffuse: blade_graphics::TextureView, out_debug: blade_graphics::TextureView, } @@ -528,6 +531,7 @@ impl Shaders { struct ShaderPipelines { fill: blade_graphics::ComputePipeline, main: blade_graphics::ComputePipeline, + spatial: blade_graphics::ComputePipeline, temporal_accum: blade_graphics::ComputePipeline, a_trous: blade_graphics::ComputePipeline, post_proc: blade_graphics::RenderPipeline, @@ -549,7 +553,7 @@ impl ShaderPipelines { compute: shader.at("main"), }) } - fn create_ray_trace( + fn create_ray_trace_main( shader: &blade_graphics::Shader, gpu: &blade_graphics::Context, ) -> blade_graphics::ComputePipeline { @@ -566,6 +570,21 @@ impl ShaderPipelines { }) } + fn create_ray_trace_spatial( + shader: &blade_graphics::Shader, + gpu: &blade_graphics::Context, + ) -> blade_graphics::ComputePipeline { + //Note: technically the spatial reuse shader could use a reduced + // set of resources, but in terms of code reuse it's much easier + // to just bind the new data as the old one. + let layout = ::layout(); + gpu.create_compute_pipeline(blade_graphics::ComputePipelineDesc { + name: "ray-trace-spatial", + data_layouts: &[&layout], + compute: shader.at("main_spatial"), + }) + } + fn create_temporal_accum( shader: &blade_graphics::Shader, gpu: &blade_graphics::Context, @@ -621,7 +640,8 @@ impl ShaderPipelines { let sh_a_trous = shader_man[shaders.a_trous].raw.as_ref().unwrap(); Ok(Self { fill: Self::create_gbuf_fill(shader_man[shaders.fill_gbuf].raw.as_ref().unwrap(), gpu), - main: Self::create_ray_trace(sh_main, gpu), + main: Self::create_ray_trace_main(sh_main, gpu), + spatial: Self::create_ray_trace_spatial(sh_main, gpu), temporal_accum: Self::create_temporal_accum(sh_a_trous, gpu), a_trous: Self::create_a_trous(sh_a_trous, gpu), post_proc: Self::create_post_proc( @@ -710,9 +730,9 @@ impl Renderer { Self { shaders, targets, - post_proc_input_index: 0, fill_pipeline: sp.fill, main_pipeline: sp.main, + spatial_pipeline: sp.spatial, post_proc_pipeline: sp.post_proc, blur: Blur { temporal_accum_pipeline: sp.temporal_accum, @@ -761,6 +781,7 @@ impl Renderer { gpu.destroy_compute_pipeline(&mut self.blur.a_trous_pipeline); gpu.destroy_compute_pipeline(&mut self.fill_pipeline); gpu.destroy_compute_pipeline(&mut self.main_pipeline); + gpu.destroy_compute_pipeline(&mut self.spatial_pipeline); gpu.destroy_render_pipeline(&mut self.post_proc_pipeline); } @@ -802,7 +823,8 @@ impl Renderer { shader.get_struct_size("StoredReservoir"), self.reservoir_size ); - self.main_pipeline = ShaderPipelines::create_ray_trace(shader, gpu); + self.main_pipeline = ShaderPipelines::create_ray_trace_main(shader, gpu); + self.spatial_pipeline = ShaderPipelines::create_ray_trace_spatial(shader, gpu); } } if self.shaders.a_trous != old.a_trous { @@ -1115,7 +1137,6 @@ impl Renderer { } self.is_frozen = config.frozen; self.targets.camera_params[self.frame_index % 2] = self.make_camera_params(camera); - self.post_proc_input_index = self.frame_index % 2; } /// Ray trace the scene. @@ -1130,7 +1151,6 @@ impl Renderer { ) { let debug = self.make_debug_params(&debug_config); let (cur, prev) = self.work_indices(); - assert_eq!(cur, self.post_proc_input_index); if let mut pass = command_encoder.compute("fill-gbuf") { let mut pc = pass.with(&self.fill_pipeline); @@ -1159,7 +1179,22 @@ impl Renderer { pc.dispatch(groups); } - if let mut pass = command_encoder.compute("ray-trace") { + let parameters = MainParams { + frame_index: self.frame_index as u32, + num_environment_samples: ray_config.num_environment_samples, + environment_importance_sampling: ray_config.environment_importance_sampling as u32, + temporal_tap: ray_config.temporal_tap as u32, + temporal_history: ray_config.temporal_history, + spatial_taps: ray_config.spatial_taps, + spatial_tap_history: ray_config.spatial_tap_history, + spatial_radius: ray_config.spatial_radius, + t_start: ray_config.t_start, + use_pairwise_mis: ray_config.pairwise_mis as u32, + defensive_mis: ray_config.defensive_mis, + use_motion_vectors: (self.frame_scene_built >= self.frame_index) as u32, + }; + + if let mut pass = command_encoder.compute("ray-trace-main") { let mut pc = pass.with(&self.main_pipeline); let groups = self.main_pipeline.get_dispatch_for(self.surface_size); pc.bind( @@ -1168,21 +1203,7 @@ impl Renderer { camera: self.targets.camera_params[cur], prev_camera: self.targets.camera_params[prev], debug, - parameters: MainParams { - frame_index: self.frame_index as u32, - num_environment_samples: ray_config.num_environment_samples, - environment_importance_sampling: ray_config.environment_importance_sampling - as u32, - temporal_tap: ray_config.temporal_tap as u32, - temporal_history: ray_config.temporal_history, - spatial_taps: ray_config.spatial_taps, - spatial_tap_history: ray_config.spatial_tap_history, - spatial_radius: ray_config.spatial_radius, - t_start: ray_config.t_start, - use_pairwise_mis: ray_config.pairwise_mis as u32, - defensive_mis: ray_config.defensive_mis, - use_motion_vectors: (self.frame_scene_built >= self.frame_index) as u32, - }, + parameters, acc_struct: self.acceleration_structure, prev_acc_struct: if self.frame_scene_built < self.frame_index || self.prev_acceleration_structure @@ -1204,9 +1225,45 @@ impl Renderer { t_prev_flat_normal: self.targets.flat_normal.views[prev], t_motion: self.targets.motion.views[0], debug_buf: self.debug.buffer_resource(), - reservoirs: self.targets.reservoir_buf[cur].into(), - prev_reservoirs: self.targets.reservoir_buf[prev].into(), - out_diffuse: self.targets.light_diffuse.views[cur], + reservoirs: self.targets.reservoir_buf[1].into(), + prev_reservoirs: self.targets.reservoir_buf[0].into(), + out_diffuse: self.targets.light_diffuse.views[1], + in_diffuse: self.targets.light_diffuse.views[0], + out_debug: self.targets.debug.views[0], + }, + ); + pc.dispatch(groups); + } + + //TODO: control by `ray_config.spatial_pass` + if let mut pass = command_encoder.compute("ray-trace-spatial") { + let mut pc = pass.with(&self.spatial_pipeline); + let groups = self.spatial_pipeline.get_dispatch_for(self.surface_size); + pc.bind( + 0, + &MainData { + camera: self.targets.camera_params[cur], + prev_camera: self.targets.camera_params[cur], + debug, + parameters, + acc_struct: self.acceleration_structure, + prev_acc_struct: self.acceleration_structure, + sampler_linear: self.samplers.linear, + sampler_nearest: self.samplers.nearest, + env_map: self.env_map.main_view, + env_weights: self.env_map.weight_view, + t_depth: self.targets.depth.views[cur], + t_prev_depth: self.targets.depth.views[cur], + t_basis: self.targets.basis.views[cur], + t_prev_basis: self.targets.basis.views[cur], + t_flat_normal: self.targets.flat_normal.views[cur], + t_prev_flat_normal: self.targets.flat_normal.views[cur], + t_motion: self.targets.motion.views[0], + debug_buf: self.debug.buffer_resource(), + reservoirs: self.targets.reservoir_buf[0].into(), + prev_reservoirs: self.targets.reservoir_buf[1].into(), + out_diffuse: self.targets.light_diffuse.views[0], + in_diffuse: self.targets.light_diffuse.views[1], out_debug: self.targets.debug.views[0], }, ); @@ -1229,7 +1286,6 @@ impl Renderer { pad: 0, }; let (cur, prev) = self.work_indices(); - let temp = 2; if denoiser_config.temporal_weight < 1.0 { let mut pass = command_encoder.compute("temporal-accum"); @@ -1244,24 +1300,23 @@ impl Renderer { camera: self.targets.camera_params[cur], prev_camera: self.targets.camera_params[prev], params, - input: self.targets.light_diffuse.views[cur], - prev_input: self.targets.light_diffuse.views[prev], + input: self.targets.light_diffuse.views[0], + prev_input: self.targets.light_diffuse.views[2], t_depth: self.targets.depth.views[cur], t_prev_depth: self.targets.depth.views[prev], t_flat_normal: self.targets.flat_normal.views[cur], t_prev_flat_normal: self.targets.flat_normal.views[prev], t_motion: self.targets.motion.views[0], - output: self.targets.light_diffuse.views[temp], + output: self.targets.light_diffuse.views[1], }, ); pc.dispatch(groups); - //Note: making `cur` contain the latest reprojection output - self.targets.light_diffuse.views.swap(cur, temp); + //Note: making `2` contain the latest reprojection output + self.targets.light_diffuse.views.swap(1, 2); } - assert_eq!(cur, self.post_proc_input_index); - let mut ping_pong = [temp, if self.is_frozen { cur } else { prev }]; for _ in 0..denoiser_config.num_passes { + self.targets.light_diffuse.views.swap(0, 1); let mut pass = command_encoder.compute("a-trous"); let mut pc = pass.with(&self.blur.a_trous_pipeline); let groups = self @@ -1272,15 +1327,14 @@ impl Renderer { 0, &ATrousData { params, - input: self.targets.light_diffuse.views[self.post_proc_input_index], + input: self.targets.light_diffuse.views + [if params.iteration == 0 { 2 } else { 1 }], t_depth: self.targets.depth.views[cur], t_flat_normal: self.targets.flat_normal.views[cur], - output: self.targets.light_diffuse.views[ping_pong[0]], + output: self.targets.light_diffuse.views[0], }, ); pc.dispatch(groups); - self.post_proc_input_index = ping_pong[0]; - ping_pong.swap(0, 1); params.iteration += 1; } } @@ -1302,7 +1356,7 @@ impl Renderer { 0, &PostProcData { t_albedo: self.targets.albedo.views[0], - light_diffuse: self.targets.light_diffuse.views[self.post_proc_input_index], + light_diffuse: self.targets.light_diffuse.views[0], t_debug: self.targets.debug.views[0], tone_map_params: ToneMapParams { enabled: 1,